Skip to content

Commit c719295

Browse files
ActhinksActhinks
authored
Merge branch 'main' into builtin-args-type-checking
2 parents 86e50b2 + 94bcd9c commit c719295

File tree

12 files changed

+161
-71
lines changed

12 files changed

+161
-71
lines changed

‎llvm/lib/MC/MCObjectStreamer.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -586,7 +586,14 @@ void MCObjectStreamer::emitCodeAlignment(Align Alignment,
586586
const MCSubtargetInfo *STI,
587587
unsigned MaxBytesToEmit) {
588588
emitValueToAlignment(Alignment, 0, 1, MaxBytesToEmit);
589-
cast<MCAlignFragment>(getCurrentFragment())->setEmitNops(true, STI);
589+
auto *F = cast<MCAlignFragment>(getCurrentFragment());
590+
F->setEmitNops(true, STI);
591+
// With RISC-V style linker relaxation, mark the section as linker-relaxable
592+
// if the alignment is larger than the minimum NOP size.
593+
unsigned Size;
594+
if (getAssembler().getBackend().shouldInsertExtraNopBytesForCodeAlign(*F,
595+
Size))
596+
getCurrentSectionOnly()->setLinkerRelaxable();
590597
}
591598

592599
void MCObjectStreamer::emitValueToOffset(const MCExpr *Offset,

‎llvm/lib/Target/RISCV/MCTargetDesc/RISCVAsmBackend.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -546,7 +546,7 @@ static uint64_t adjustFixupValue(const MCFixup &Fixup, uint64_t Value,
546546

547547
bool RISCVAsmBackend::isPCRelFixupResolved(const MCSymbol *SymA,
548548
const MCFragment &F) {
549-
// If the section does not contain linker-relaxable instructions, PC-relative
549+
// If the section does not contain linker-relaxable fragments, PC-relative
550550
// fixups can be resolved.
551551
if (!F.getParent()->isLinkerRelaxable())
552552
return true;

‎llvm/lib/Target/RISCV/RISCVInstrInfo.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4189,6 +4189,24 @@ bool RISCVInstrInfo::simplifyInstruction(MachineInstr &MI) const {
41894189
return true;
41904190
}
41914191
break;
4192+
case RISCV::BLTU:
4193+
// bltu zero, rs, imm => bne rs, zero, imm
4194+
if (MI.getOperand(0).getReg() == RISCV::X0) {
4195+
MachineOperand MO0 = MI.getOperand(0);
4196+
MI.removeOperand(0);
4197+
MI.insert(MI.operands_begin() + 1, {MO0});
4198+
MI.setDesc(get(RISCV::BNE));
4199+
}
4200+
break;
4201+
case RISCV::BGEU:
4202+
// bgeu zero, rs, imm => beq rs, zero, imm
4203+
if (MI.getOperand(0).getReg() == RISCV::X0) {
4204+
MachineOperand MO0 = MI.getOperand(0);
4205+
MI.removeOperand(0);
4206+
MI.insert(MI.operands_begin() + 1, {MO0});
4207+
MI.setDesc(get(RISCV::BEQ));
4208+
}
4209+
break;
41924210
}
41934211
return false;
41944212
}

‎llvm/test/CodeGen/RISCV/machine-copyprop-simplifyinstruction.mir

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -742,3 +742,39 @@ body: |
742742
renamable $x10 = MAXU renamable $x11, renamable $x11
743743
PseudoRET implicit $x10
744744
...
745+
---
746+
name: bltu
747+
body: |
748+
; CHECK-LABEL: name: bltu
749+
; CHECK: bb.0:
750+
; CHECK-NEXT: successors: %bb.1(0x80000000)
751+
; CHECK-NEXT: {{ $}}
752+
; CHECK-NEXT: renamable $x11 = COPY $x12
753+
; CHECK-NEXT: BNE $x12, $x0, %bb.1
754+
; CHECK-NEXT: {{ $}}
755+
; CHECK-NEXT: bb.1:
756+
; CHECK-NEXT: PseudoRET
757+
bb.0:
758+
renamable $x11 = COPY $x12
759+
BLTU $x0, renamable $x11, %bb.1
760+
bb.1:
761+
PseudoRET
762+
...
763+
---
764+
name: bgeu
765+
body: |
766+
; CHECK-LABEL: name: bgeu
767+
; CHECK: bb.0:
768+
; CHECK-NEXT: successors: %bb.1(0x80000000)
769+
; CHECK-NEXT: {{ $}}
770+
; CHECK-NEXT: renamable $x11 = COPY $x12
771+
; CHECK-NEXT: BEQ $x12, $x0, %bb.1
772+
; CHECK-NEXT: {{ $}}
773+
; CHECK-NEXT: bb.1:
774+
; CHECK-NEXT: PseudoRET
775+
bb.0:
776+
renamable $x11 = COPY $x12
777+
BGEU $x0, renamable $x11, %bb.1
778+
bb.1:
779+
PseudoRET
780+
...

‎llvm/test/DebugInfo/RISCV/relax-debug-frame.ll

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -4,12 +4,12 @@
44
; RUN: | FileCheck -check-prefix=RELAX-DWARFDUMP %s
55
;
66
; RELAX: Section ({{.*}}) .rela.eh_frame {
7-
; REALX-NEXT: 0x1C R_RISCV_32_PCREL .L0 0x0
8-
; REALX-NEXT: 0x30 R_RISCV_32_PCREL .L0 0x0
9-
; REALX-NEXT: 0x48 R_RISCV_32_PCREL .L0 0x0
10-
; REALX-NEXT: 0x4C R_RISCV_ADD32 .L0 0x0
11-
; REALX-NEXT: 0x4C R_RISCV_SUB32 .L0 0x0
12-
; REALX-NEXT: 0x57 R_RISCV_SET6 .L0 0x0
7+
; RELAX-NEXT: 0x1C R_RISCV_32_PCREL .L0 0x0
8+
; RELAX-NEXT: 0x30 R_RISCV_32_PCREL .L0 0x0
9+
; RELAX-NEXT: 0x48 R_RISCV_32_PCREL .L0 0x0
10+
; RELAX-NEXT: 0x4C R_RISCV_ADD32 .L0 0x0
11+
; RELAX-NEXT: 0x4C R_RISCV_SUB32 .L0 0x0
12+
; RELAX-NEXT: 0x57 R_RISCV_SET6 .L0 0x0
1313
; RELAX-NEXT-EMPTY:
1414

1515
; RELAX-DWARFDUMP-NOT: error: failed to compute relocation

‎llvm/test/MC/RISCV/align.s

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -134,3 +134,23 @@ data2:
134134
.option norvc
135135
.balign 4
136136
add a0, a0, a1
137+
138+
## Branches crossing the linker-relaxable R_RISCV_ALIGN need relocations.
139+
# RELAX-RELOC: .rela.text3 {
140+
# RELAX-RELOC-NEXT: 0x4 R_RISCV_BRANCH .Ltmp[[#]] 0x0
141+
# RELAX-RELOC-NEXT: 0x8 R_RISCV_ALIGN - 0x4
142+
# RELAX-RELOC-NEXT: 0xC R_RISCV_BRANCH .Ltmp[[#]] 0x0
143+
# RELAX-RELOC-NEXT: }
144+
# C-OR-ZCA-EXT-RELAX-RELOC: .rela.text3 {
145+
# C-OR-ZCA-EXT-RELAX-RELOC-NEXT: 0x4 R_RISCV_BRANCH .Ltmp[[#]] 0x0
146+
# C-OR-ZCA-EXT-RELAX-RELOC-NEXT: 0x8 R_RISCV_ALIGN - 0x4
147+
# C-OR-ZCA-EXT-RELAX-RELOC-NEXT: 0xC R_RISCV_BRANCH .Ltmp[[#]] 0x0
148+
# C-OR-ZCA-EXT-RELAX-RELOC-NEXT: }
149+
.section .text3, "ax"
150+
bnez t1, 1f
151+
bnez t2, 2f
152+
1:
153+
.p2align 3
154+
2:
155+
bnez t1, 1b
156+
bnez t1, 2b

‎llvm/test/MC/RISCV/cfi-advance.s

Lines changed: 19 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -13,16 +13,26 @@
1313
# RELAX: R_RISCV_CALL_PLT
1414
# RELAX-NEXT: R_RISCV_RELAX
1515
# RELAX-EMPTY:
16-
# CHECK-NEXT: Relocation section '.rela.eh_frame' at offset {{.*}} contains 1 entries:
17-
# CHECK: Offset Info Type Sym. Value Symbol's Name + Addend
18-
# CHECK-NEXT: 0000001c 00000139 R_RISCV_32_PCREL 00000000 .L0 + 0
16+
# NORELAX-NEXT: Relocation section '.rela.eh_frame' at offset {{.*}} contains 1 entries:
17+
# NORELAX: Offset Info Type Sym. Value Symbol's Name + Addend
18+
# NORELAX-NEXT: 0000001c 00000139 R_RISCV_32_PCREL 00000000 .L0 + 0
19+
# RELAX-NEXT: Relocation section '.rela.eh_frame' at offset {{.*}} contains 5 entries:
20+
# RELAX: Offset Info Type Sym. Value Symbol's Name + Addend
21+
# RELAX-NEXT: 0000001c 00000139 R_RISCV_32_PCREL 00000000 .L0 + 0
22+
# RELAX-NEXT: 00000020 00000c23 R_RISCV_ADD32 0001017a .L0 + 0
23+
# RELAX-NEXT: 00000020 00000127 R_RISCV_SUB32 00000000 .L0 + 0
24+
# RELAX-NEXT: 00000035 00000b35 R_RISCV_SET6 00010176 .L0 + 0
25+
# RELAX-NEXT: 00000035 00000934 R_RISCV_SUB6 0001016e .L0 + 0
1926
# CHECK-EMPTY:
20-
# CHECK: Symbol table '.symtab' contains 13 entries:
21-
# CHECK-NEXT: Num: Value Size Type Bind Vis Ndx Name
22-
# CHECK-NEXT: 0: 00000000 0 NOTYPE LOCAL DEFAULT UND
23-
# CHECK-NEXT: 1: 00000000 0 NOTYPE LOCAL DEFAULT 2 .L0 {{$}}
24-
# CHECK: 3: 00000004 0 NOTYPE LOCAL DEFAULT 2 .L0{{$}}
25-
# CHECK-NOT: .L0
27+
# NORELAX: Symbol table '.symtab' contains 13 entries:
28+
# RELAX: Symbol table '.symtab' contains 16 entries:
29+
# RELAX-NEXT: Num: Value Size Type Bind Vis Ndx Name
30+
# RELAX-NEXT: 0: 00000000 0 NOTYPE LOCAL DEFAULT UND
31+
# RELAX-NEXT: 1: 00000000 0 NOTYPE LOCAL DEFAULT 2 .L0 {{$}}
32+
# RELAX: 3: 00000004 0 NOTYPE LOCAL DEFAULT 2 .L0{{$}}
33+
# RELAX: 9: 0001016e 0 NOTYPE LOCAL DEFAULT 2 .L0 {{$}}
34+
# RELAX: 11: 00010176 0 NOTYPE LOCAL DEFAULT 2 .L0 {{$}}
35+
# RELAX: 12: 0001017a 0 NOTYPE LOCAL DEFAULT 2 .L0 {{$}}
2636

2737
# CHECK-DWARFDUMP: DW_CFA_advance_loc1: 104
2838
# CHECK-DWARFDUMP-NEXT: DW_CFA_def_cfa_offset: +8

‎llvm/tools/llvm-jitlink/llvm-jitlink.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,6 @@
1717
#include "llvm/Config/llvm-config.h" // for LLVM_ON_UNIX, LLVM_ENABLE_THREADS
1818
#include "llvm/ExecutionEngine/Orc/AbsoluteSymbols.h"
1919
#include "llvm/ExecutionEngine/Orc/COFFPlatform.h"
20-
#include "llvm/ExecutionEngine/Orc/COFFVCRuntimeSupport.h"
2120
#include "llvm/ExecutionEngine/Orc/DebugObjectManagerPlugin.h"
2221
#include "llvm/ExecutionEngine/Orc/Debugging/DebugInfoSupport.h"
2322
#include "llvm/ExecutionEngine/Orc/Debugging/DebuggerSupportPlugin.h"

‎mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -3640,36 +3640,38 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
36403640
}
36413641

36423642
//===----------------------------------------------------------------------===//
3643-
// NVVM dot.accumulate.4way Op
3643+
// NVVM dot.accumulate Ops
36443644
//===----------------------------------------------------------------------===//
36453645

3646-
def DotAccumulate4WayS8 : I32EnumAttrCase<"S8", 1, "s8">;
3647-
def DotAccumulate4WayU8 : I32EnumAttrCase<"U8", 0, "u8">;
3646+
def DotAccumulateUnsigned : I32EnumAttrCase<"UNSIGNED", 0, "unsigned">;
3647+
def DotAccumulateSigned : I32EnumAttrCase<"SIGNED", 1, "signed">;
36483648

3649-
def DotAccumulate4WayType : I32EnumAttr<"DotAccumulate4WayType",
3650-
"NVVM DotAccumulate4WayType",
3651-
[DotAccumulate4WayS8, DotAccumulate4WayU8]> {
3649+
def DotAccumulateType : I32EnumAttr<"DotAccumulateType",
3650+
"NVVM DotAccumulateType",
3651+
[DotAccumulateSigned, DotAccumulateUnsigned]> {
36523652
let cppNamespace = "::mlir::NVVM";
36533653
let genSpecializedAttr = 0;
36543654
}
36553655

3656-
def DotAccumulate4WayTypeAttr : EnumAttr<NVVM_Dialect, DotAccumulate4WayType, "dot_accumulate_4way_type"> {
3656+
def DotAccumulateTypeAttr : EnumAttr<NVVM_Dialect, DotAccumulateType, "dot_accumulate_type"> {
36573657
let assemblyFormat = "`<` $value `>`";
36583658
}
36593659

36603660
def NVVM_DotAccumulate4WayOp : NVVM_Op<"dot.accumulate.4way"> {
3661-
let summary = "Four-way byte dot product-accumulate instruction.";
3661+
let summary = "Four-way byte dot product-accumulate instruction";
36623662
let description = [{
36633663
Performs a four-way byte dot-product which is accumulated in a 32-bit
36643664
result.
36653665
Operand `a` and `b` are vectors of 4 bytes between which the dot product is
36663666
computed.
3667+
36673668
The `a_type` and `b_type` attributes specify the type of the elements in `a`
36683669
and `b` respectively.
3669-
If `a_type` or `b_type` is `s8`, then the elements in the corresponding
3670+
If `a_type` or `b_type` is `signed`, then the elements in the corresponding
36703671
vector are sign-extended to 32-bit before the dot product is computed.
3671-
If `a_type` or `b_type` is `u8`, then the elements in the corresponding
3672-
vector are zero-extended to 32-bit instead.
3672+
If `a_type` or `b_type` is `unsigned`, then the elements in the
3673+
corresponding vector are zero-extended to 32-bit instead.
3674+
36733675
Operand `c` is a 32-bit integer to which the result is accumulated. It is
36743676
treated as holding a signed integer if any of `a_type` or `b_type` is `s8`.
36753677

@@ -3678,9 +3680,9 @@ def NVVM_DotAccumulate4WayOp : NVVM_Op<"dot.accumulate.4way"> {
36783680

36793681
let arguments = (ins
36803682
VectorOfLengthAndType<[4], [I8]>:$a,
3681-
DotAccumulate4WayTypeAttr:$a_type,
3683+
DotAccumulateTypeAttr:$a_type,
36823684
VectorOfLengthAndType<[4], [I8]>:$b,
3683-
DotAccumulate4WayTypeAttr:$b_type,
3685+
DotAccumulateTypeAttr:$b_type,
36843686
I32:$c
36853687
);
36863688

@@ -3689,17 +3691,15 @@ def NVVM_DotAccumulate4WayOp : NVVM_Op<"dot.accumulate.4way"> {
36893691
let assemblyFormat = "$a $a_type `,` $b $b_type `,` $c attr-dict `:` type($a) `,` type($b)";
36903692

36913693
let extraClassDeclaration = [{
3692-
static llvm::Intrinsic::ID
3693-
getIntrinsicID(NVVM::DotAccumulate4WayType a_type,
3694-
NVVM::DotAccumulate4WayType b_type);
3695-
llvm::Value* getPackedArg(llvm::Value* arg, llvm::IRBuilderBase& builder);
3694+
static mlir::NVVM::IDArgPair
3695+
getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
3696+
llvm::IRBuilderBase &builder);
36963697
}];
36973698

36983699
string llvmBuilder = [{
3699-
llvm::Intrinsic::ID id = NVVM::DotAccumulate4WayOp::getIntrinsicID($a_type, $b_type);
3700-
llvm::Value* argA = op.getPackedArg($a, builder);
3701-
llvm::Value* argB = op.getPackedArg($b, builder);
3702-
$res = createIntrinsicCall(builder, id, {argA, argB, $c});
3700+
auto [id, args] = NVVM::DotAccumulate4WayOp::getIntrinsicIDAndArgs(
3701+
*op, moduleTranslation, builder);
3702+
$res = createIntrinsicCall(builder, id, args);
37033703
}];
37043704
}
37053705

‎mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

Lines changed: 25 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1205,13 +1205,6 @@ LogicalResult NVVM::VoteSyncOp::verify() {
12051205
return success();
12061206
}
12071207

1208-
llvm::Value *
1209-
NVVM::DotAccumulate4WayOp::getPackedArg(llvm::Value *arg,
1210-
llvm::IRBuilderBase &builder) {
1211-
return builder.CreateBitCast(arg,
1212-
llvm::Type::getInt32Ty(builder.getContext()));
1213-
}
1214-
12151208
/// Packs the given `field` into the `result`.
12161209
/// The `result` is 64-bits and each `field` can be 32-bits or narrower.
12171210
static llvm::Value *
@@ -1692,24 +1685,31 @@ static void nvvmInferResultRanges(Operation *op, Value result,
16921685
}
16931686
}
16941687

1695-
llvm::Intrinsic::ID
1696-
DotAccumulate4WayOp::getIntrinsicID(NVVM::DotAccumulate4WayType a_type,
1697-
NVVM::DotAccumulate4WayType b_type) {
1698-
bool is_a_siext = a_type == NVVM::DotAccumulate4WayType::S8;
1699-
bool is_b_siext = b_type == NVVM::DotAccumulate4WayType::S8;
1700-
unsigned type = (is_a_siext << 1) | is_b_siext;
1701-
switch (type) {
1702-
case 0:
1703-
return llvm::Intrinsic::nvvm_idp4a_u_u;
1704-
case 1:
1705-
return llvm::Intrinsic::nvvm_idp4a_u_s;
1706-
case 2:
1707-
return llvm::Intrinsic::nvvm_idp4a_s_u;
1708-
case 3:
1709-
return llvm::Intrinsic::nvvm_idp4a_s_s;
1710-
default:
1711-
llvm_unreachable("Invalid DP4a type");
1712-
}
1688+
static llvm::Value *getAsPackedI32(llvm::Value *arg,
1689+
llvm::IRBuilderBase &builder) {
1690+
return builder.CreateBitCast(arg,
1691+
llvm::Type::getInt32Ty(builder.getContext()));
1692+
}
1693+
1694+
NVVM::IDArgPair DotAccumulate4WayOp::getIntrinsicIDAndArgs(
1695+
Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
1696+
auto curOp = cast<NVVM::DotAccumulate4WayOp>(op);
1697+
1698+
llvm::SmallVector<llvm::Value *> args;
1699+
args.push_back(getAsPackedI32(mt.lookupValue(curOp.getA()), builder));
1700+
args.push_back(getAsPackedI32(mt.lookupValue(curOp.getB()), builder));
1701+
args.push_back(mt.lookupValue(curOp.getC()));
1702+
1703+
bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
1704+
bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
1705+
unsigned type = (isASigned << 1) | isBSigned;
1706+
const llvm::Intrinsic::ID ids[] = {
1707+
llvm::Intrinsic::nvvm_idp4a_u_u,
1708+
llvm::Intrinsic::nvvm_idp4a_u_s,
1709+
llvm::Intrinsic::nvvm_idp4a_s_u,
1710+
llvm::Intrinsic::nvvm_idp4a_s_s,
1711+
};
1712+
return {ids[type], args};
17131713
}
17141714

17151715
//===----------------------------------------------------------------------===//

‎mlir/test/Dialect/LLVMIR/nvvm.mlir

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -579,11 +579,11 @@ func.func @st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size: i64)
579579
}
580580

581581
// CHECK-LABEL: @dot_accumulate_4way
582-
func.func @dot_accumulate_4way(%a: i32, %a_vec: vector<4xi8>, %b: i32, %b_vec: vector<4xi8>, %c: i32) {
582+
func.func @dot_accumulate_4way(%a_vec: vector<4xi8>, %b_vec: vector<4xi8>, %c: i32) {
583583
// CHECK: nvvm.dot.accumulate.4way %{{.*}}, %{{.*}}, %{{.*}} : vector<4xi8>, vector<4xi8>
584-
%1 = nvvm.dot.accumulate.4way %a_vec <u8>, %b_vec <u8>, %c: vector<4xi8>, vector<4xi8>
584+
%1 = nvvm.dot.accumulate.4way %a_vec <unsigned>, %b_vec <unsigned>, %c: vector<4xi8>, vector<4xi8>
585585
// CHECK: nvvm.dot.accumulate.4way %{{.*}}, %{{.*}}, %{{.*}} : vector<4xi8>, vector<4xi8>
586-
%3 = nvvm.dot.accumulate.4way %a_vec <s8>, %b_vec <s8>, %c: vector<4xi8>, vector<4xi8>
586+
%3 = nvvm.dot.accumulate.4way %a_vec <signed>, %b_vec <signed>, %c: vector<4xi8>, vector<4xi8>
587587
return
588588
}
589589

‎mlir/test/Target/LLVMIR/nvvmir.mlir

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -851,18 +851,18 @@ llvm.func @nvvm_dot_accumulate_4way(%a: vector<4xi8>, %b: vector<4xi8>, %c: i32)
851851
// CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
852852
// CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
853853
// CHECK: call i32 @llvm.nvvm.idp4a.u.u(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})
854-
%0 = nvvm.dot.accumulate.4way %a <u8>, %b <u8>, %c: vector<4xi8>, vector<4xi8>
854+
%0 = nvvm.dot.accumulate.4way %a <unsigned>, %b <unsigned>, %c: vector<4xi8>, vector<4xi8>
855855
// CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
856856
// CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
857857
// CHECK: call i32 @llvm.nvvm.idp4a.s.u(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})
858-
%1 = nvvm.dot.accumulate.4way %a <s8>, %b <u8>, %c: vector<4xi8>, vector<4xi8>
858+
%1 = nvvm.dot.accumulate.4way %a <signed>, %b <unsigned>, %c: vector<4xi8>, vector<4xi8>
859859
// CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
860860
// CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
861861
// CHECK: call i32 @llvm.nvvm.idp4a.u.s(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})
862-
%2 = nvvm.dot.accumulate.4way %a <u8>, %b <s8>, %c: vector<4xi8>, vector<4xi8>
862+
%2 = nvvm.dot.accumulate.4way %a <unsigned>, %b <signed>, %c: vector<4xi8>, vector<4xi8>
863863
// CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
864864
// CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32
865865
// CHECK: call i32 @llvm.nvvm.idp4a.s.s(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})
866-
%3 = nvvm.dot.accumulate.4way %a <s8>, %b <s8>, %c: vector<4xi8>, vector<4xi8>
866+
%3 = nvvm.dot.accumulate.4way %a <signed>, %b <signed>, %c: vector<4xi8>, vector<4xi8>
867867
llvm.return
868868
}

0 commit comments

Comments
 (0)