[llvm] [mlir] [SDAG][NVPTX] Enable custom legalization of v1 type for Intrinsic results and operands (PR #203237)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 11 03:39:43 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-llvm
Author: Pradeep Kumar (schwarzschild-radius)
<details>
<summary>Changes</summary>
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
---
Patch is 51.77 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/203237.diff
10 Files Affected:
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+3-10)
- (modified) llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp (+8)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+14-7)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+8-3)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-ld.ll (+6-6)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-st.ll (+16-16)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+7-11)
- (modified) mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir (+3-3)
- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir (+12-12)
- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir (+22-22)
``````````diff
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_16...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/203237
More information about the llvm-commits
mailing list