[llvm] [mlir] [SDAG][NVPTX] Enable custom legalization of v1 type for Intrinsic results and operands (PR #203237)
Pradeep Kumar via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 11 03:39:02 PDT 2026
https://github.com/schwarzschild-radius created https://github.com/llvm/llvm-project/pull/203237
This commit enables custom legalization for v1 types for intrinsic results and operand by calling CustomLowerNode which gives the target a chance to custom handle the node before scalarizing
To demonstrate, I have used tcgen05.ld/st intrinsics which previously used a scalar i32 type while all wider variants used vector types
Assisted by: Claude Code
>From cf52eb67caad981260b6f150b40d05dfb32f01c8 Mon Sep 17 00:00:00 2001
From: Pradeep Kumar <pradeepku at nvidia.com>
Date: Mon, 23 Mar 2026 09:25:43 +0000
Subject: [PATCH] [SDAG] Enable custom legalization of v1 types
This commit enables custom legalization for v1 types by calling CustomLowerNode
which gives the target a chance to custom handle the node before scalarizing
To demonstrate, I have used tcgen05.ld/st intrinsics which previously used a
scalar i32 type while all wider variants used vector types
Assisted by: Claude Code
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 13 ++----
.../SelectionDAG/LegalizeVectorTypes.cpp | 8 ++++
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 21 ++++++---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 11 +++--
llvm/test/CodeGen/NVPTX/tcgen05-ld.ll | 12 ++---
llvm/test/CodeGen/NVPTX/tcgen05-st.ll | 32 +++++++-------
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 18 +++-----
.../Dialect/LLVMIR/nvvm_check_target_sm.mlir | 6 +--
mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir | 24 +++++-----
mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir | 44 +++++++++----------
10 files changed, 99 insertions(+), 90 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 361746c853160..bd53223296d69 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1145,7 +1145,8 @@ class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
[OpType, llvm_i32_ty, llvm_i32_ty]);
}
-class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num, string ElemType = "i32"> {
+class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num,
+ string ElemType = "i32"> {
int shift = !cond(!eq(Shape, "16x128b"): 1,
!eq(Shape, "16x256b"): 2,
true : 0);
@@ -1153,15 +1154,7 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num, string ElemType = "i3
int veclen = !shl(1, !add(Num, shift));
int valid = !le(veclen, 128);
- LLVMType type = !cond(!eq(veclen, 1): LLVMType<!cast<ValueType>(ElemType)>,
- !eq(veclen, 2): LLVMType<!cast<ValueType>("v"#2#ElemType)>,
- !eq(veclen, 4): LLVMType<!cast<ValueType>("v"#4#ElemType)>,
- !eq(veclen, 8): LLVMType<!cast<ValueType>("v"#8#ElemType)>,
- !eq(veclen, 16): LLVMType<!cast<ValueType>("v"#16#ElemType)>,
- !eq(veclen, 32): LLVMType<!cast<ValueType>("v"#32#ElemType)>,
- !eq(veclen, 64): LLVMType<!cast<ValueType>("v"#64#ElemType)>,
- !eq(veclen, 128): LLVMType<!cast<ValueType>("v"#128#ElemType)>,
- true : llvm_void_ty);
+ LLVMType type = LLVMType<!cast<ValueType>("v"#veclen#ElemType)>;
}
class NVVM_TCGEN05_MMA_BASE<string Space, bit IsSparse, string Kind = ""> {
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index 5350be412176d..91922ea64555f 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -43,6 +43,10 @@ void DAGTypeLegalizer::ScalarizeVectorResult(SDNode *N, unsigned ResNo) {
N->dump(&DAG));
SDValue R = SDValue();
+ // See if the target wants to custom expand this node.
+ if (CustomLowerNode(N, N->getValueType(ResNo), true))
+ return;
+
switch (N->getOpcode()) {
default:
#ifndef NDEBUG
@@ -839,6 +843,10 @@ bool DAGTypeLegalizer::ScalarizeVectorOperand(SDNode *N, unsigned OpNo) {
N->dump(&DAG));
SDValue Res = SDValue();
+ // See if the target wants to custom scalarize this node.
+ if (CustomLowerNode(N, N->getOperand(OpNo).getValueType(), false))
+ return false;
+
switch (N->getOpcode()) {
default:
#ifndef NDEBUG
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 17d9f857312d6..48ba3542888c6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1123,16 +1123,17 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
// Custom lowering for tcgen05.ld vector operands
setOperationAction(ISD::INTRINSIC_W_CHAIN,
- {MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
- MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::v2f32,
- MVT::v4f32, MVT::v8f32, MVT::v16f32, MVT::v32f32,
- MVT::v64f32, MVT::v128f32},
+ {MVT::v1i32, MVT::v2i32, MVT::v4i32, MVT::v8i32,
+ MVT::v16i32, MVT::v32i32, MVT::v64i32, MVT::v128i32,
+ MVT::v2f32, MVT::v4f32, MVT::v8f32, MVT::v16f32,
+ MVT::v32f32, MVT::v64f32, MVT::v128f32},
Custom);
// Custom lowering for tcgen05.st vector operands
setOperationAction(ISD::INTRINSIC_VOID,
- {MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
- MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::Other},
+ {MVT::v1i32, MVT::v2i32, MVT::v4i32, MVT::v8i32,
+ MVT::v16i32, MVT::v32i32, MVT::v64i32, MVT::v128i32,
+ MVT::Other},
Custom);
// Enable custom lowering for the following:
@@ -2817,6 +2818,7 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
switch (IntrinNo) {
default:
break;
+ case Intrinsic::nvvm_tcgen05_st_16x64b_x1:
case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
case Intrinsic::nvvm_tcgen05_st_16x64b_x8:
@@ -2836,6 +2838,7 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
+ case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
@@ -2845,6 +2848,7 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
return lowerTcgen05St(Op, DAG);
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
@@ -5376,7 +5380,7 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1: {
Info.opc = ISD::INTRINSIC_VOID;
- Info.memVT = MVT::i32;
+ Info.memVT = MVT::v1i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
@@ -7268,12 +7272,14 @@ static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG,
return;
}
+ case Intrinsic::nvvm_tcgen05_ld_16x64b_x1:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x4:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x8:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x16:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x32:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x64:
case Intrinsic::nvvm_tcgen05_ld_16x64b_x128:
+ case Intrinsic::nvvm_tcgen05_ld_32x32b_x1:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x4:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x8:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x16:
@@ -7298,6 +7304,7 @@ static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG,
}
return;
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1:
case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4:
case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8:
case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16:
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 75e7e4cbbf0eb..2ff421e5a6386 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5769,7 +5769,10 @@ class TCGEN05_LDST_REGINFO<int Veclen> {
list<NVPTXRegClass> regs = !listsplat(B32, Veclen);
// generate list of regnames for load/store operands
list<string> reg_names = !foreach(x, !range(0, Veclen), "r" # x);
- string regstring = "{{" # !interleave(!foreach(n, !range(0, Veclen), "$r" # n), ", ") # "}}";
+ string regstring = "{{" #
+ !interleave(
+ !foreach(n, !range(0, Veclen), "$r" # n), ", ") #
+ "}}";
dag Ins = !dag(ins, regs, reg_names);
dag Outs = !dag(outs, regs, reg_names);
}
@@ -5785,7 +5788,8 @@ class TCGEN05_LD_INST<string Shape, int Num, bit Pack> :
NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.veclen>;
let InOperandList = !con((ins B32:$taddr),
- !if(!eq(Shape, "16x32bx2"), (ins i64imm:$offset), (ins)));
+ !if(!eq(Shape, "16x32bx2"),
+ (ins i64imm:$offset), (ins)));
let OutOperandList = Info.Outs;
let AsmString = "tcgen05.ld.sync.aligned"
# "." # Shape
@@ -5809,7 +5813,8 @@ class TCGEN05_ST_INST<string Shape, int Num, bit Unpack> :
NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.veclen>;
let InOperandList = !con((ins B32:$taddr),
- !if(!eq(Shape, "16x32bx2"), (ins i64imm:$offset), (ins)),
+ !if(!eq(Shape, "16x32bx2"),
+ (ins i64imm:$offset), (ins)),
Info.Ins);
let OutOperandList = (outs);
let AsmString = "tcgen05.st.sync.aligned"
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-ld.ll b/llvm/test/CodeGen/NVPTX/tcgen05-ld.ll
index 22eb7298133bb..33eb39ac7701e 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-ld.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-ld.ll
@@ -27,7 +27,7 @@ define void @nvvm_tcgen05_ld_16x64b(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x64b.x64.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1];
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x64b.x128.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1];
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) %taddr, i1 0)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) %taddr, i1 0)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x2(ptr addrspace(6) %taddr, i1 0)
@@ -62,7 +62,7 @@ define void @nvvm_tcgen05_ld_16x64b_pack(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x64b.x64.pack::16b.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1];
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x64b.x128.pack::16b.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1];
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) %taddr, i1 1)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) %taddr, i1 1)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x2(ptr addrspace(6) %taddr, i1 1)
@@ -219,7 +219,7 @@ define void @nvvm_tcgen05_ld_32x32b(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.32x32b.x64.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1];
; CHECK-NEXT: tcgen05.ld.sync.aligned.32x32b.x128.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1];
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) %taddr, i1 0)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) %taddr, i1 0)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x2(ptr addrspace(6) %taddr, i1 0)
@@ -253,7 +253,7 @@ define void @nvvm_tcgen05_ld_32x32b_pack(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.32x32b.x64.pack::16b.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1];
; CHECK-NEXT: tcgen05.ld.sync.aligned.32x32b.x128.pack::16b.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1];
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) %taddr, i1 1)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) %taddr, i1 1)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x2(ptr addrspace(6) %taddr, i1 1)
@@ -288,7 +288,7 @@ define void @nvvm_tcgen05_ld_16x32bx2(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x32bx2.x64.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1], 2;
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x32bx2.x128.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1], 2;
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i1 0)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i1 0)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, i1 0)
@@ -322,7 +322,7 @@ define void @nvvm_tcgen05_ld_16x32bx2_pack(ptr addrspace(6) %taddr) {
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x32bx2.x64.pack::16b.b32 {%r65, %r66, %r67, %r68, %r69, %r70, %r71, %r72, %r73, %r74, %r75, %r76, %r77, %r78, %r79, %r80, %r81, %r82, %r83, %r84, %r85, %r86, %r87, %r88, %r89, %r90, %r91, %r92, %r93, %r94, %r95, %r96, %r97, %r98, %r99, %r100, %r101, %r102, %r103, %r104, %r105, %r106, %r107, %r108, %r109, %r110, %r111, %r112, %r113, %r114, %r115, %r116, %r117, %r118, %r119, %r120, %r121, %r122, %r123, %r124, %r125, %r126, %r127, %r128}, [%r1], 2;
; CHECK-NEXT: tcgen05.ld.sync.aligned.16x32bx2.x128.pack::16b.b32 {%r129, %r130, %r131, %r132, %r133, %r134, %r135, %r136, %r137, %r138, %r139, %r140, %r141, %r142, %r143, %r144, %r145, %r146, %r147, %r148, %r149, %r150, %r151, %r152, %r153, %r154, %r155, %r156, %r157, %r158, %r159, %r160, %r161, %r162, %r163, %r164, %r165, %r166, %r167, %r168, %r169, %r170, %r171, %r172, %r173, %r174, %r175, %r176, %r177, %r178, %r179, %r180, %r181, %r182, %r183, %r184, %r185, %r186, %r187, %r188, %r189, %r190, %r191, %r192, %r193, %r194, %r195, %r196, %r197, %r198, %r199, %r200, %r201, %r202, %r203, %r204, %r205, %r206, %r207, %r208, %r209, %r210, %r211, %r212, %r213, %r214, %r215, %r216, %r217, %r218, %r219, %r220, %r221, %r222, %r223, %r224, %r225, %r226, %r227, %r228, %r229, %r230, %r231, %r232, %r233, %r234, %r235, %r236, %r237, %r238, %r239, %r240, %r241, %r242, %r243, %r244, %r245, %r246, %r247, %r248, %r249, %r250, %r251, %r252, %r253, %r254, %r255, %r256}, [%r1], 2;
; CHECK-NEXT: ret;
- tail call i32 @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i1 1)
+ tail call <1 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i1 1)
tail call <2 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, i1 1)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-st.ll b/llvm/test/CodeGen/NVPTX/tcgen05-st.ll
index ccf6541d01973..952e1cf39b6fc 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-st.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-st.ll
@@ -11,7 +11,7 @@
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
; CHECK-LABEL: nvvm_tcgen05_st_16x64b
-define void @nvvm_tcgen05_st_16x64b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x64b(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x64b(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -92,7 +92,7 @@ define void @nvvm_tcgen05_st_16x64b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x64b_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.16x64b.x128.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 0)
+ tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, <1 x i32> %stv1, i1 0)
tail call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 0)
@@ -111,7 +111,7 @@ define void @nvvm_tcgen05_st_16x64b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32
}
; CHECK-LABEL: nvvm_tcgen05_st_16x64b_unpack
-define void @nvvm_tcgen05_st_16x64b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x64b_unpack(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x64b_unpack(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -192,7 +192,7 @@ define void @nvvm_tcgen05_st_16x64b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x64b_unpack_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.16x64b.x128.unpack::16b.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 1)
+ tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, <1 x i32> %stv1, i1 1)
tail call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 1)
@@ -211,7 +211,7 @@ define void @nvvm_tcgen05_st_16x64b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <
}
; CHECK-LABEL: nvvm_tcgen05_st_16x128b
-define void @nvvm_tcgen05_st_16x128b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x128b(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x128b(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<256>;
@@ -307,7 +307,7 @@ define void @nvvm_tcgen05_st_16x128b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i3
}
; CHECK-LABEL: nvvm_tcgen05_st_16x128b_unpack
-define void @nvvm_tcgen05_st_16x128b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x128b_unpack(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x128b_unpack(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<256>;
@@ -403,7 +403,7 @@ define void @nvvm_tcgen05_st_16x128b_unpack(ptr addrspace(6) %taddr, i32 %stv1,
}
; CHECK-LABEL: nvvm_tcgen05_st_16x256b
-define void @nvvm_tcgen05_st_16x256b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x256b(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x256b(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<254>;
@@ -495,7 +495,7 @@ define void @nvvm_tcgen05_st_16x256b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i3
}
; CHECK-LABEL: nvvm_tcgen05_st_16x256b_unpack
-define void @nvvm_tcgen05_st_16x256b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x256b_unpack(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x256b_unpack(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<254>;
@@ -587,7 +587,7 @@ define void @nvvm_tcgen05_st_16x256b_unpack(ptr addrspace(6) %taddr, i32 %stv1,
}
; CHECK-LABEL: nvvm_tcgen05_st_32x32b
-define void @nvvm_tcgen05_st_32x32b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_32x32b(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_32x32b(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -668,7 +668,7 @@ define void @nvvm_tcgen05_st_32x32b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_32x32b_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.32x32b.x128.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 0)
+ tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, <1 x i32> %stv1, i1 0)
tail call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 0)
@@ -687,7 +687,7 @@ define void @nvvm_tcgen05_st_32x32b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32
}
; CHECK-LABEL: nvvm_tcgen05_st_32x32b_unpack
-define void @nvvm_tcgen05_st_32x32b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_32x32b_unpack(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_32x32b_unpack(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -768,7 +768,7 @@ define void @nvvm_tcgen05_st_32x32b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_32x32b_unpack_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.32x32b.x128.unpack::16b.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 1)
+ tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, <1 x i32> %stv1, i1 1)
tail call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 1)
@@ -787,7 +787,7 @@ define void @nvvm_tcgen05_st_32x32b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <
}
; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2
-define void @nvvm_tcgen05_st_16x32bx2(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x32bx2(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -868,7 +868,7 @@ define void @nvvm_tcgen05_st_16x32bx2(ptr addrspace(6) %taddr, i32 %stv1, <2 x i
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x32bx2_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.16x32bx2.x128.b32 [%r1], 2, {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i32 %stv1, i1 0)
+ tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, <1 x i32> %stv1, i1 0)
tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, <2 x i32> %stv2, i1 0)
@@ -887,7 +887,7 @@ define void @nvvm_tcgen05_st_16x32bx2(ptr addrspace(6) %taddr, i32 %stv1, <2 x i
}
; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2_unpack
-define void @nvvm_tcgen05_st_16x32bx2_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
+define void @nvvm_tcgen05_st_16x32bx2_unpack(ptr addrspace(6) %taddr, <1 x i32> %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {
; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2_unpack(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<257>;
@@ -968,7 +968,7 @@ define void @nvvm_tcgen05_st_16x32bx2_unpack(ptr addrspace(6) %taddr, i32 %stv1,
; CHECK-NEXT: ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8];
; CHECK-NEXT: tcgen05.st.sync.aligned.16x32bx2.x128.unpack::16b.b32 [%r1], 2, {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};
; CHECK-NEXT: ret;
- tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i32 %stv1, i1 1)
+ tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, <1 x i32> %stv1, i1 1)
tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, <2 x i32> %stv2, i1 1)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index dabb519aa1395..cfaa8df158f56 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5540,8 +5540,8 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMf<[100, 101, 110]>]>
Optional<I64>:$offset
);
- let results = (outs AnyTypeOf<[I32, VectorOfLengthAndType<
- [2, 4, 8, 16, 32, 64, 128], [I32]>]>:$res);
+ let results = (outs VectorOfLengthAndType<
+ [1, 2, 4, 8, 16, 32, 64, 128], [I32]>:$res);
let assemblyFormat = [{
$tmemAddr (`,` $offset^)? (`pack` $pack^)? attr-dict `:` type($res)
@@ -5600,11 +5600,9 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMf<[100, 101, 110]>]>
llvm::LLVMContext &Context = moduleTranslation.getLLVMContext();
auto Pack = llvm::ConstantInt::get(Context, llvm::APInt(1, $pack));
- unsigned num = $_resultType->isVectorTy()
- ? llvm::cast<llvm::VectorType>($_resultType)
+ unsigned num = llvm::cast<llvm::VectorType>($_resultType)
->getElementCount()
- .getFixedValue()
- : 1;
+ .getFixedValue();
auto ID = getTcgen05LdIntrinsicID($shape, num);
if (ID == llvm::Intrinsic::not_intrinsic)
@@ -5719,8 +5717,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMf<[100, 101, 110]>]>
Tcgen05LdStShapeAttr:$shape,
// Arguments
LLVM_PointerTensor:$tmemAddr,
- AnyTypeOf<[I32, VectorOfLengthAndType<
- [2, 4, 8, 16, 32, 64, 128], [I32]>]>:$val,
+ VectorOfLengthAndType<[1, 2, 4, 8, 16, 32, 64, 128], [I32]>:$val,
Optional<I64>:$offset
);
@@ -5779,10 +5776,9 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMf<[100, 101, 110]>]>
auto Unpack = llvm::ConstantInt::get(Context, llvm::APInt(1, $unpack));
auto valTy = $val->getType();
- uint32_t num = valTy->isVectorTy() ? llvm::cast<llvm::VectorType>(valTy)
+ uint32_t num = llvm::cast<llvm::VectorType>(valTy)
->getElementCount()
- .getFixedValue()
- : 1;
+ .getFixedValue();
auto ID = getTcgen05StIntrinsicID($shape, num);
if (ID == llvm::Intrinsic::not_intrinsic)
diff --git a/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir b/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
index ff90ad47ba410..34f7f5b87e72b 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
@@ -83,7 +83,7 @@ gpu.module @tcgen05_cp_sm90a [#nvvm.target<chip = "sm_90a">] {
gpu.module @tcgen05_ld_sm90 [#nvvm.target<chip = "sm_90">] {
func.func @tcgen05_ld_sm90(%taddr: !llvm.ptr<6>) {
// expected-error @below {{'nvvm.tcgen05.ld' op is not supported on sm_90}}
- %0 = nvvm.tcgen05.ld %taddr {shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : i32
+ %0 = nvvm.tcgen05.ld %taddr {shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<1 x i32>
return
}
}
@@ -91,9 +91,9 @@ gpu.module @tcgen05_ld_sm90 [#nvvm.target<chip = "sm_90">] {
// -----
gpu.module @tcgen05_st_sm120f [#nvvm.target<chip = "sm_120f">] {
- func.func @tcgen05_st_sm120f(%taddr: !llvm.ptr<6>, %val: i32) {
+ func.func @tcgen05_st_sm120f(%taddr: !llvm.ptr<6>, %val: vector<1 x i32>) {
// expected-error @below {{'nvvm.tcgen05.st' op is not supported on sm_120f}}
- nvvm.tcgen05.st %taddr, %val {shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : i32
+ nvvm.tcgen05.st %taddr, %val {shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<1 x i32>
return
}
}
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir
index b1266b0e8151d..bcd342e2e6e94 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir
@@ -3,8 +3,8 @@
// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b
llvm.func @nvvm_tcgen05_ld_16x64b(%tmemAddr : !llvm.ptr<6>) {
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 false)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 false)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, i1 false)
%ldv2 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<2 x i32>
@@ -33,8 +33,8 @@ llvm.func @nvvm_tcgen05_ld_16x64b(%tmemAddr : !llvm.ptr<6>) {
// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b_pack
llvm.func @nvvm_tcgen05_ld_16x64b_pack(%tmemAddr : !llvm.ptr<6>) {
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 true)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 true)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, i1 true)
%ldv2 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>} : vector<2 x i32>
@@ -165,8 +165,8 @@ llvm.func @nvvm_tcgen05_ld_16x256b_pack(%tmemAddr : !llvm.ptr<6>) {
// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b
llvm.func @nvvm_tcgen05_ld_32x32b(%tmemAddr : !llvm.ptr<6>) {
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 false)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 false)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, i1 false)
%ldv2 = nvvm.tcgen05.ld %tmemAddr { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : vector<2 x i32>
@@ -195,8 +195,8 @@ llvm.func @nvvm_tcgen05_ld_32x32b(%tmemAddr : !llvm.ptr<6>) {
// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b_pack
llvm.func @nvvm_tcgen05_ld_32x32b_pack(%tmemAddr : !llvm.ptr<6>) {
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 true)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i1 true)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, i1 true)
%ldv2 = nvvm.tcgen05.ld %tmemAddr pack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>} : vector<2 x i32>
@@ -227,8 +227,8 @@ llvm.func @nvvm_tcgen05_ld_16x32bx2(%tmemAddr : !llvm.ptr<6>) {
%halfSplitOffset = llvm.mlir.constant(2:i64) : i64
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 false)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 false)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 false)
%ldv2 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : vector<2 x i32>
@@ -259,8 +259,8 @@ llvm.func @nvvm_tcgen05_ld_16x32bx2_pack(%tmemAddr : !llvm.ptr<6>) {
%halfSplitOffset = llvm.mlir.constant(2:i64) : i64
-// CHECK: call i32 @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 true)
- %ldv1 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : i32
+// CHECK: call <1 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 true)
+ %ldv1 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : vector<1 x i32>
// CHECK: call <2 x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, i1 true)
%ldv2 = nvvm.tcgen05.ld %tmemAddr, %halfSplitOffset pack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>} : vector<2 x i32>
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir
index 119746133625d..123c6308794a6 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir
@@ -3,7 +3,7 @@
// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b
llvm.func @nvvm_tcgen05_ld_16x64b(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -12,8 +12,8 @@ llvm.func @nvvm_tcgen05_ld_16x64b(
%stv64 : vector<64xi32>,
%stv128 : vector<128xi32>) {
-// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 false)
- nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, <1 x i32> {{%[0-9]+}}, i1 false)
+ nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 false)
nvvm.tcgen05.st %tmemAddr, %stv2 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=2:i32 } : vector<2xi32>
@@ -42,7 +42,7 @@ llvm.func @nvvm_tcgen05_ld_16x64b(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b_pack
llvm.func @nvvm_tcgen05_ld_16x64b_pack(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -51,8 +51,8 @@ llvm.func @nvvm_tcgen05_ld_16x64b_pack(
%stv64 : vector<64xi32>,
%stv128 : vector<128xi32>) {
-// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 true)
- nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, <1 x i32> {{%[0-9]+}}, i1 true)
+ nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 true)
nvvm.tcgen05.st %tmemAddr, %stv2 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=2:i32 } : vector<2xi32>
@@ -81,7 +81,7 @@ llvm.func @nvvm_tcgen05_ld_16x64b_pack(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x128b
llvm.func @nvvm_tcgen05_ld_16x128b(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -117,7 +117,7 @@ llvm.func @nvvm_tcgen05_ld_16x128b(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x128b_pack
llvm.func @nvvm_tcgen05_ld_16x128b_pack(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -153,7 +153,7 @@ llvm.func @nvvm_tcgen05_ld_16x128b_pack(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x256b
llvm.func @nvvm_tcgen05_ld_16x256b(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -186,7 +186,7 @@ llvm.func @nvvm_tcgen05_ld_16x256b(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x256b_pack
llvm.func @nvvm_tcgen05_ld_16x256b_pack(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -219,7 +219,7 @@ llvm.func @nvvm_tcgen05_ld_16x256b_pack(
// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b
llvm.func @nvvm_tcgen05_ld_32x32b(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -228,8 +228,8 @@ llvm.func @nvvm_tcgen05_ld_32x32b(
%stv64 : vector<64xi32>,
%stv128 : vector<128xi32>) {
-// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 false)
- nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, <1 x i32> {{%[0-9]+}}, i1 false)
+ nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 false)
nvvm.tcgen05.st %tmemAddr, %stv2 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=2:i32 } : vector<2xi32>
@@ -258,7 +258,7 @@ llvm.func @nvvm_tcgen05_ld_32x32b(
// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b_pack
llvm.func @nvvm_tcgen05_ld_32x32b_pack(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -267,8 +267,8 @@ llvm.func @nvvm_tcgen05_ld_32x32b_pack(
%stv64 : vector<64xi32>,
%stv128 : vector<128xi32>) {
-// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 true)
- nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, <1 x i32> {{%[0-9]+}}, i1 true)
+ nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 true)
nvvm.tcgen05.st %tmemAddr, %stv2 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=2:i32 } : vector<2xi32>
@@ -297,7 +297,7 @@ llvm.func @nvvm_tcgen05_ld_32x32b_pack(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x32bx2
llvm.func @nvvm_tcgen05_ld_16x32bx2(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -308,8 +308,8 @@ llvm.func @nvvm_tcgen05_ld_16x32bx2(
%offset = llvm.mlir.constant(2:i64) : i64
-// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i32 {{%[0-9]+}}, i1 false)
- nvvm.tcgen05.st %tmemAddr, %stv1, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, <1 x i32> {{%[0-9]+}}, i1 false)
+ nvvm.tcgen05.st %tmemAddr, %stv1, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, <2 x i32> {{%[0-9]+}}, i1 false)
nvvm.tcgen05.st %tmemAddr, %stv2, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=2:i32 } : vector<2xi32>
@@ -338,7 +338,7 @@ llvm.func @nvvm_tcgen05_ld_16x32bx2(
// CHECK-LABEL: @nvvm_tcgen05_ld_16x32bx2_pack
llvm.func @nvvm_tcgen05_ld_16x32bx2_pack(
%tmemAddr : !llvm.ptr<6>,
- %stv1 : i32,
+ %stv1 : vector<1 x i32>,
%stv2 : vector<2xi32>,
%stv4 : vector<4xi32>,
%stv8 : vector<8xi32>,
@@ -349,8 +349,8 @@ llvm.func @nvvm_tcgen05_ld_16x32bx2_pack(
%offset = llvm.mlir.constant(2:i64) : i64
-// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i32 {{%[0-9]+}}, i1 true)
- nvvm.tcgen05.st %tmemAddr, %stv1, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : i32
+// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, <1 x i32> {{%[0-9]+}}, i1 true)
+ nvvm.tcgen05.st %tmemAddr, %stv1, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : vector<1 x i32>
// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, <2 x i32> {{%[0-9]+}}, i1 true)
nvvm.tcgen05.st %tmemAddr, %stv2, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=2:i32 } : vector<2xi32>
More information about the llvm-commits
mailing list