Skip to content

Commit ba7d479

Browse files
durga4githubdurga4github
authored
[MLIR][NVVM] [NFC] Rename Tcgen05GroupKind to CTAGroupKind (#156448)
...as the cta_group::1/2 are used in non-tcgen05 Ops like TMA Loads also. Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
1 parent 9c3961f commit ba7d479

File tree

6 files changed

+53
-53
lines changed

6 files changed

+53
-53
lines changed

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

Lines changed: 25 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -2368,6 +2368,23 @@ def TMAStoreModeAttr : EnumAttr<NVVM_Dialect, TMAStoreMode, "tma_store_mode"> {
23682368
let assemblyFormat = "`<` $value `>`";
23692369
}
23702370

2371+
// Num CTAs in a group participating in the TMA/MMA operations.
2372+
// This corresponds to the "cta_group::1", "cta_group::2"
2373+
// modifiers in the PTX instructions.
2374+
def CTAGroup_1 : I32EnumAttrCase<"CTA_1", 0, "cta_1">;
2375+
def CTAGroup_2 : I32EnumAttrCase<"CTA_2", 1, "cta_2">;
2376+
2377+
def CTAGroupKind : I32EnumAttr<"CTAGroupKind",
2378+
"NVVM CTA group kind",
2379+
[CTAGroup_1, CTAGroup_2]> {
2380+
let genSpecializedAttr = 0;
2381+
let cppNamespace = "::mlir::NVVM";
2382+
}
2383+
def CTAGroupKindAttr :
2384+
EnumAttr<NVVM_Dialect, CTAGroupKind, "cta_group"> {
2385+
let assemblyFormat = "`<` $value `>`";
2386+
}
2387+
23712388
def NVVM_CpAsyncBulkCommitGroupOp : NVVM_Op<"cp.async.bulk.commit.group">,
23722389
Arguments<(ins )> {
23732390
let assemblyFormat = "attr-dict";
@@ -3333,23 +3350,6 @@ def NVVM_Breakpoint : NVVM_Op<"breakpoint"> {
33333350
//===----------------------------------------------------------------------===//
33343351
// NVVM TCGEN05 Ops
33353352
//===----------------------------------------------------------------------===//
3336-
// Num CTAs in a group participating in the TCGEN05 operation.
3337-
// This corresponds to the "cta_group::1", "cta_group::2"
3338-
// modifiers in the PTX instructions.
3339-
def Tcgen05GroupCTA_1 : I32EnumAttrCase<"CTA_1", 0, "cta_1">;
3340-
def Tcgen05GroupCTA_2 : I32EnumAttrCase<"CTA_2", 1, "cta_2">;
3341-
3342-
def Tcgen05GroupKind : I32EnumAttr<"Tcgen05GroupKind",
3343-
"NVVM Tcgen05 group kind",
3344-
[Tcgen05GroupCTA_1, Tcgen05GroupCTA_2]> {
3345-
let genSpecializedAttr = 0;
3346-
let cppNamespace = "::mlir::NVVM";
3347-
}
3348-
def Tcgen05GroupKindAttr :
3349-
EnumAttr<NVVM_Dialect, Tcgen05GroupKind, "tcgen05_group"> {
3350-
let assemblyFormat = "`<` $value `>`";
3351-
}
3352-
33533353
def Tcgen05FenceBefore : I32EnumAttrCase<"BEFORE_THREAD_SYNC", 0, "before">;
33543354
def Tcgen05FenceAfter : I32EnumAttrCase<"AFTER_THREAD_SYNC", 1, "after">;
33553355
def Tcgen05FenceKind : I32EnumAttr<"Tcgen05FenceKind", "NVVM Tcgen05 fence kind",
@@ -3387,7 +3387,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMa<[100, 101]>]
33873387
let arguments = (ins
33883388
AnyTypeOf<[LLVM_AnyPointer, LLVM_PointerShared]>:$addr,
33893389
I32:$nCols,
3390-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group);
3390+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group);
33913391

33923392
let assemblyFormat = "$addr `,` $nCols attr-dict `:` type(operands)";
33933393

@@ -3415,7 +3415,7 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMa<[100, 10
34153415
}];
34163416

34173417
let arguments = (ins LLVM_PointerTensor:$taddr, I32:$nCols,
3418-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group);
3418+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group);
34193419

34203420
let assemblyFormat = "$taddr `,` $nCols attr-dict `:` type(operands)";
34213421

@@ -3443,12 +3443,12 @@ def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_perm
34433443
}];
34443444

34453445
let arguments = (ins
3446-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group);
3446+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group);
34473447

34483448
let assemblyFormat = "attr-dict";
34493449

34503450
string llvmBuilder = [{
3451-
auto id = ($group == NVVM::Tcgen05GroupKind::CTA_1) ?
3451+
auto id = ($group == NVVM::CTAGroupKind::CTA_1) ?
34523452
llvm::Intrinsic::nvvm_tcgen05_relinq_alloc_permit_cg1 :
34533453
llvm::Intrinsic::nvvm_tcgen05_relinq_alloc_permit_cg2;
34543454
createIntrinsicCall(builder, id);
@@ -3516,7 +3516,7 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMa<[100, 101]
35163516
let arguments = (ins
35173517
AnyTypeOf<[LLVM_AnyPointer, LLVM_PointerShared]>:$addr,
35183518
Optional<I16>:$multicastMask,
3519-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group);
3519+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group);
35203520

35213521
let assemblyFormat = [{
35223522
$addr (`,` `multicast_mask` `=` $multicastMask^)?
@@ -3549,12 +3549,12 @@ def NVVM_Tcgen05ShiftOp : NVVM_Op<"tcgen05.shift", [NVVMRequiresSMa<[100, 101, 1
35493549
}];
35503550

35513551
let arguments = (ins LLVM_PointerTensor:$taddr,
3552-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group);
3552+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group);
35533553

35543554
let assemblyFormat = "$taddr attr-dict `:` type(operands)";
35553555

35563556
string llvmBuilder = [{
3557-
auto id = ($group == NVVM::Tcgen05GroupKind::CTA_1) ?
3557+
auto id = ($group == NVVM::CTAGroupKind::CTA_1) ?
35583558
llvm::Intrinsic::nvvm_tcgen05_shift_down_cg1 :
35593559
llvm::Intrinsic::nvvm_tcgen05_shift_down_cg2;
35603560
createIntrinsicCall(builder, id, {$taddr});
@@ -3626,7 +3626,7 @@ def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMa<[100, 101]>]> {
36263626

36273627
let arguments = (ins
36283628
Tcgen05CpShapeAttr:$shape,
3629-
DefaultValuedAttr<Tcgen05GroupKindAttr, "Tcgen05GroupKind::CTA_1">:$group,
3629+
DefaultValuedAttr<CTAGroupKindAttr, "CTAGroupKind::CTA_1">:$group,
36303630
DefaultValuedAttr<Tcgen05CpMulticastAttr, "Tcgen05CpMulticast::NONE">:$multicast,
36313631
OptionalAttr<Tcgen05CpSrcFormatAttr>:$srcFormat,
36323632
LLVM_PointerTensor:$taddr,

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1797,7 +1797,7 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
17971797
unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
17981798
.getAddressSpace();
17991799
bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
1800-
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
1800+
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
18011801

18021802
llvm::Intrinsic::ID id;
18031803
if (isShared) {
@@ -1819,7 +1819,7 @@ llvm::Intrinsic::ID Tcgen05DeallocOp::getIntrinsicIDAndArgs(
18191819
Operation &op, LLVM::ModuleTranslation &mt,
18201820
llvm::SmallVector<llvm::Value *> &args) {
18211821
auto curOp = cast<NVVM::Tcgen05DeallocOp>(op);
1822-
auto id = (curOp.getGroup() == Tcgen05GroupKind::CTA_1)
1822+
auto id = (curOp.getGroup() == CTAGroupKind::CTA_1)
18231823
? llvm::Intrinsic::nvvm_tcgen05_dealloc_cg1
18241824
: llvm::Intrinsic::nvvm_tcgen05_dealloc_cg2;
18251825

@@ -1847,7 +1847,7 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
18471847
.getAddressSpace();
18481848
bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
18491849
bool hasMulticast = static_cast<bool>(curOp.getMulticastMask());
1850-
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
1850+
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
18511851

18521852
llvm::Intrinsic::ID id =
18531853
is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
@@ -1879,7 +1879,7 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
18791879

18801880
llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(Operation &op) {
18811881
auto curOp = cast<NVVM::Tcgen05CpOp>(op);
1882-
bool is2CTA = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
1882+
bool is2CTA = curOp.getGroup() == CTAGroupKind::CTA_2;
18831883
auto srcFmt = curOp.getSrcFormat();
18841884
auto mc = curOp.getMulticast();
18851885

‎mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@ llvm.func @llvm_nvvm_tcgen05_alloc(%addr : !llvm.ptr, %ncols : i32) {
66
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr, i32
77

88
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %{{.*}}, i32 %{{.*}})
9-
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr, i32
9+
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i32
1010
llvm.return
1111
}
1212

@@ -16,7 +16,7 @@ llvm.func @llvm_nvvm_tcgen05_alloc_shared(%addr : !llvm.ptr<3>, %ncols : i32) {
1616
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr<3>, i32
1717

1818
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
19-
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr<3>, i32
19+
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i32
2020
llvm.return
2121
}
2222

@@ -26,7 +26,7 @@ llvm.func @llvm_nvvm_tcgen05_dealloc(%addr : !llvm.ptr<6>, %ncols : i32) {
2626
nvvm.tcgen05.dealloc %addr, %ncols : !llvm.ptr<6>, i32
2727

2828
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %{{.*}}, i32 %{{.*}})
29-
nvvm.tcgen05.dealloc %addr, %ncols {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr<6>, i32
29+
nvvm.tcgen05.dealloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<6>, i32
3030
llvm.return
3131
}
3232

@@ -36,6 +36,6 @@ llvm.func @llvm_nvvm_tcgen05_relinquish_alloc_permit() {
3636
nvvm.tcgen05.relinquish_alloc_permit
3737

3838
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()
39-
nvvm.tcgen05.relinquish_alloc_permit {group = #nvvm.tcgen05_group<cta_2>}
39+
nvvm.tcgen05.relinquish_alloc_permit {group = #nvvm.cta_group<cta_2>}
4040
llvm.return
4141
}

‎mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6,13 +6,13 @@ llvm.func @llvm_nvvm_tcgen05_commit_generic(%barrier : !llvm.ptr, %cta_mask : i1
66
nvvm.tcgen05.commit %barrier : !llvm.ptr
77

88
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg2(ptr %{{.*}})
9-
nvvm.tcgen05.commit %barrier {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr
9+
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr
1010

1111
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %{{.*}}, i16 %{{.*}})
1212
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr, i16
1313

1414
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %{{.*}}, i16 %{{.*}})
15-
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr, i16
15+
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i16
1616
llvm.return
1717
}
1818

@@ -22,12 +22,12 @@ llvm.func @llvm_nvvm_tcgen05_commit_shared(%barrier : !llvm.ptr<3>, %cta_mask :
2222
nvvm.tcgen05.commit %barrier : !llvm.ptr<3>
2323

2424
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %{{.*}})
25-
nvvm.tcgen05.commit %barrier {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr<3>
25+
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>
2626

2727
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
2828
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr<3>, i16
2929

3030
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
31-
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr<3>, i16
31+
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i16
3232
llvm.return
3333
}

‎mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir‎

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -6,18 +6,18 @@ llvm.func @nvvm_tcgen05_cp_128x256b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
66
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>}
77

88
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
9-
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>, group = #nvvm.tcgen05_group<cta_2>}
9+
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>, group = #nvvm.cta_group<cta_2>}
1010

1111
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
1212
nvvm.tcgen05.cp %taddr, %smem_desc {
1313
shape = #nvvm.tcgen05_cp_shape<shape_128x256b>,
14-
group = #nvvm.tcgen05_group<cta_2>,
14+
group = #nvvm.cta_group<cta_2>,
1515
srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>
1616
}
1717
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
1818
nvvm.tcgen05.cp %taddr, %smem_desc {
1919
shape = #nvvm.tcgen05_cp_shape<shape_128x256b>,
20-
group = #nvvm.tcgen05_group<cta_2>,
20+
group = #nvvm.cta_group<cta_2>,
2121
srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>
2222
}
2323
llvm.return
@@ -29,18 +29,18 @@ llvm.func @nvvm_tcgen05_cp_4x256b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
2929
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_4x256b>}
3030

3131
// CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
32-
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_4x256b>, group = #nvvm.tcgen05_group<cta_2>}
32+
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_4x256b>, group = #nvvm.cta_group<cta_2>}
3333

3434
// CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
3535
nvvm.tcgen05.cp %taddr, %smem_desc {
3636
shape = #nvvm.tcgen05_cp_shape<shape_4x256b>,
37-
group = #nvvm.tcgen05_group<cta_2>,
37+
group = #nvvm.cta_group<cta_2>,
3838
srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>
3939
}
4040
// CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
4141
nvvm.tcgen05.cp %taddr, %smem_desc {
4242
shape = #nvvm.tcgen05_cp_shape<shape_4x256b>,
43-
group = #nvvm.tcgen05_group<cta_2>,
43+
group = #nvvm.cta_group<cta_2>,
4444
srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>
4545
}
4646
llvm.return
@@ -52,18 +52,18 @@ llvm.func @nvvm_tcgen05_cp_128x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
5252
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x128b>}
5353

5454
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
55-
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x128b>, group = #nvvm.tcgen05_group<cta_2>}
55+
nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x128b>, group = #nvvm.cta_group<cta_2>}
5656

5757
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
5858
nvvm.tcgen05.cp %taddr, %smem_desc {
5959
shape = #nvvm.tcgen05_cp_shape<shape_128x128b>,
60-
group = #nvvm.tcgen05_group<cta_2>,
60+
group = #nvvm.cta_group<cta_2>,
6161
srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>
6262
}
6363
// CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
6464
nvvm.tcgen05.cp %taddr, %smem_desc {
6565
shape = #nvvm.tcgen05_cp_shape<shape_128x128b>,
66-
group = #nvvm.tcgen05_group<cta_2>,
66+
group = #nvvm.cta_group<cta_2>,
6767
srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>
6868
}
6969
llvm.return
@@ -80,21 +80,21 @@ llvm.func @nvvm_tcgen05_cp_64x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
8080
// CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
8181
nvvm.tcgen05.cp %taddr, %smem_desc {
8282
shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,
83-
group = #nvvm.tcgen05_group<cta_2>,
83+
group = #nvvm.cta_group<cta_2>,
8484
multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>
8585
}
8686

8787
// CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
8888
nvvm.tcgen05.cp %taddr, %smem_desc {
8989
shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,
90-
group = #nvvm.tcgen05_group<cta_1>,
90+
group = #nvvm.cta_group<cta_1>,
9191
multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>,
9292
srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>
9393
}
9494
// CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
9595
nvvm.tcgen05.cp %taddr, %smem_desc {
9696
shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,
97-
group = #nvvm.tcgen05_group<cta_2>,
97+
group = #nvvm.cta_group<cta_2>,
9898
multicast = #nvvm.tcgen05_cp_multicast<warpx2_01_23>,
9999
srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>
100100
}
@@ -113,21 +113,21 @@ llvm.func @nvvm_tcgen05_cp_32x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
113113
// CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
114114
nvvm.tcgen05.cp %taddr, %smem_desc {
115115
shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,
116-
group = #nvvm.tcgen05_group<cta_2>,
116+
group = #nvvm.cta_group<cta_2>,
117117
multicast = #nvvm.tcgen05_cp_multicast<warpx4>
118118
}
119119

120120
// CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
121121
nvvm.tcgen05.cp %taddr, %smem_desc {
122122
shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,
123-
group = #nvvm.tcgen05_group<cta_2>,
123+
group = #nvvm.cta_group<cta_2>,
124124
multicast = #nvvm.tcgen05_cp_multicast<warpx4>,
125125
srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>
126126
}
127127
// CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})
128128
nvvm.tcgen05.cp %taddr, %smem_desc {
129129
shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,
130-
group = #nvvm.tcgen05_group<cta_1>,
130+
group = #nvvm.cta_group<cta_1>,
131131
multicast = #nvvm.tcgen05_cp_multicast<warpx4>,
132132
srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>
133133
}

‎mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,6 @@ llvm.func @llvm_nvvm_tcgen05_shift(%taddr : !llvm.ptr<6>) {
66
nvvm.tcgen05.shift %taddr : !llvm.ptr<6>
77

88
// CHECK: call void @llvm.nvvm.tcgen05.shift.down.cg2(ptr addrspace(6) %{{.*}})
9-
nvvm.tcgen05.shift %taddr {group = #nvvm.tcgen05_group<cta_2>} : !llvm.ptr<6>
9+
nvvm.tcgen05.shift %taddr {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<6>
1010
llvm.return
1111
}

0 commit comments

Comments
 (0)