diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index fc7abc50dce01..d96f94cf4889f 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -1145,7 +1145,8 @@ class SHFL_INFO { [OpType, llvm_i32_ty, llvm_i32_ty]); } -class NVVM_TCGEN05_LDST_ACCESS_SIZE { +class NVVM_TCGEN05_LDST_ACCESS_SIZE { int shift = !cond(!eq(Shape, "16x128b"): 1, !eq(Shape, "16x256b"): 2, true : 0); @@ -1153,15 +1154,7 @@ class NVVM_TCGEN05_LDST_ACCESS_SIZE, - !eq(veclen, 4): LLVMType("v"#4#ElemType)>, - !eq(veclen, 8): LLVMType("v"#8#ElemType)>, - !eq(veclen, 16): LLVMType("v"#16#ElemType)>, - !eq(veclen, 32): LLVMType("v"#32#ElemType)>, - !eq(veclen, 64): LLVMType("v"#64#ElemType)>, - !eq(veclen, 128): LLVMType("v"#128#ElemType)>, - true : llvm_void_ty); + LLVMType type = LLVMType("v"#veclen#ElemType)>; } class NVVM_TCGEN05_MMA_BASE { diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp index 547ee11db45e2..3a511b5d022c9 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 @@ -860,6 +864,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 fa0b9fe5f7051..bacd857b19e8b 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -1124,19 +1124,19 @@ 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 and the st.async // i128 (.b128) operand. MVT::i8 is needed for the st.async.{sys,gpu} b8 // variant. setOperationAction(ISD::INTRINSIC_VOID, - {MVT::i8, MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32, - MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::i128, - MVT::Other}, + {MVT::i8, MVT::v1i32, MVT::v2i32, MVT::v4i32, MVT::v8i32, + MVT::v16i32, MVT::v32i32, MVT::v64i32, MVT::v128i32, + MVT::i128, MVT::Other}, Custom); // Enable custom lowering for the following: @@ -2796,6 +2796,8 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) { case Intrinsic::nvvm_st_async_gpu: case Intrinsic::nvvm_st_async_mmio_sys: return lowerStAsyncRelease(Op, DAG); + + 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: @@ -2815,6 +2817,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: @@ -2824,6 +2827,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: @@ -5357,7 +5361,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; @@ -7249,12 +7253,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: @@ -7279,6 +7285,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 3819b3143f447..5af5ae2fe9e52 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -5769,7 +5769,10 @@ class TCGEN05_LDST_REGINFO { list regs = !listsplat(B32, Veclen); // generate list of regnames for load/store operands list 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 : NVVM_TCGEN05_LDST_ACCESS_SIZE.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 : NVVM_TCGEN05_LDST_ACCESS_SIZE.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 40f7f15b694cb..89950d95e647f 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -5538,8 +5538,8 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMf<[100, 101, 110]>]> Optional:$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) @@ -5598,11 +5598,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($_resultType) + unsigned num = llvm::cast($_resultType) ->getElementCount() - .getFixedValue() - : 1; + .getFixedValue(); auto ID = getTcgen05LdIntrinsicID($shape, num); if (ID == llvm::Intrinsic::not_intrinsic) @@ -5717,8 +5715,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:$offset ); @@ -5777,10 +5774,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(valTy) + uint32_t num = llvm::cast(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] { gpu.module @tcgen05_ld_sm90 [#nvvm.target] { 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} : i32 + %0 = nvvm.tcgen05.ld %taddr {shape = #nvvm.tcgen05_ldst_shape} : vector<1 x i32> return } } @@ -91,9 +91,9 @@ gpu.module @tcgen05_ld_sm90 [#nvvm.target] { // ----- gpu.module @tcgen05_st_sm120f [#nvvm.target] { - 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} : i32 + nvvm.tcgen05.st %taddr, %val {shape = #nvvm.tcgen05_ldst_shape} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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} : 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, 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, num=2:i32 } : vector<2xi32>