Skip to content

Commit 78c6e07

Browse files
Wolfram70Wolfram70
committed
fix warnings
1 parent 84b8e76 commit 78c6e07

File tree

2 files changed

+36
-29
lines changed

2 files changed

+36
-29
lines changed

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

Lines changed: 24 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -416,7 +416,7 @@ def NVVM_ReduxOp :
416416
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
417417
}];
418418
string llvmBuilder = [{
419-
auto [id, args, _] =
419+
auto [id, args, types] =
420420
NVVM::ReduxOp::getIIDAndArgsWithTypes(
421421
*op, moduleTranslation, builder);
422422
$res = createIntrinsicCall(builder, id, args);
@@ -890,7 +890,7 @@ def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
890890
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
891891
}];
892892
let llvmBuilder = [{
893-
auto [intId, args, _] =
893+
auto [intId, args, types] =
894894
NVVM::FenceProxyAcquireOp::getIIDAndArgsWithTypes(
895895
*op, moduleTranslation, builder);
896896
createIntrinsicCall(builder, intId, args);
@@ -924,10 +924,9 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
924924
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
925925
}];
926926
let llvmBuilder = [{
927-
auto [intId, _, _] =
928-
NVVM::FenceProxyReleaseOp::getIIDAndArgsWithTypes(
927+
auto [intId, args, types] = NVVM::FenceProxyReleaseOp::getIIDAndArgsWithTypes(
929928
*op, moduleTranslation, builder);
930-
createIntrinsicCall(builder, intId);
929+
createIntrinsicCall(builder, intId, args);
931930
}];
932931

933932
let hasVerifier = 1;
@@ -1012,7 +1011,7 @@ def NVVM_ShflOp :
10121011
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
10131012
}];
10141013
string llvmBuilder = [{
1015-
auto [intId, args, _] = NVVM::ShflOp::getIIDAndArgsWithTypes(
1014+
auto [intId, args, types] = NVVM::ShflOp::getIIDAndArgsWithTypes(
10161015
*op, moduleTranslation, builder);
10171016
$res = createIntrinsicCall(builder, intId, args);
10181017
}];
@@ -1066,7 +1065,7 @@ def NVVM_VoteSyncOp
10661065
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
10671066
}];
10681067
string llvmBuilder = [{
1069-
auto [intId, args, _] =
1068+
auto [intId, args, types] =
10701069
NVVM::VoteSyncOp::getIIDAndArgsWithTypes(
10711070
*op, moduleTranslation, builder);
10721071
$res = createIntrinsicCall(builder, intId, args);
@@ -1145,7 +1144,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
11451144
llvm::IRBuilderBase &builder);
11461145
}];
11471146
string llvmBuilder = [{
1148-
auto [id, args, _] = NVVM::CpAsyncOp::getIIDAndArgsWithTypes(
1147+
auto [id, args, types] = NVVM::CpAsyncOp::getIIDAndArgsWithTypes(
11491148
*op, moduleTranslation, builder);
11501149
createIntrinsicCall(builder, id, args);
11511150
}];
@@ -2608,7 +2607,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
26082607
let hasVerifier = 1;
26092608

26102609
string llvmBuilder = [{
2611-
auto [id, args, _] = NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2610+
auto [id, args, types] = NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
26122611
*op, moduleTranslation, builder);
26132612
createIntrinsicCall(builder, id, args);
26142613
}];
@@ -2686,7 +2685,7 @@ def NVVM_PrefetchOp : NVVM_Op<"prefetch",
26862685
}
26872686
}];
26882687
let llvmBuilder = [{
2689-
auto [id, args, _] = NVVM::PrefetchOp::getIIDAndArgsWithTypes(*op,
2688+
auto [id, args, types] = NVVM::PrefetchOp::getIIDAndArgsWithTypes(*op,
26902689
moduleTranslation, builder);
26912690

26922691
if(op.getTensormap())
@@ -2734,7 +2733,7 @@ def NVVM_CpAsyncBulkPrefetchOp : NVVM_Op<"cp.async.bulk.prefetch"> {
27342733
}];
27352734

27362735
string llvmBuilder = [{
2737-
auto [id, args, _] = NVVM::CpAsyncBulkPrefetchOp::getIIDAndArgsWithTypes(
2736+
auto [id, args, types] = NVVM::CpAsyncBulkPrefetchOp::getIIDAndArgsWithTypes(
27382737
*op, moduleTranslation, builder);
27392738
createIntrinsicCall(builder, id, args);
27402739
}];
@@ -2777,7 +2776,7 @@ def NVVM_CpAsyncBulkTensorPrefetchOp :
27772776
let hasVerifier = 1;
27782777

27792778
string llvmBuilder = [{
2780-
auto [id, args, _] = NVVM::CpAsyncBulkTensorPrefetchOp::getIIDAndArgsWithTypes(
2779+
auto [id, args, types] = NVVM::CpAsyncBulkTensorPrefetchOp::getIIDAndArgsWithTypes(
27812780
*op, moduleTranslation, builder);
27822781
createIntrinsicCall(builder, id, args);
27832782
}];
@@ -2846,7 +2845,7 @@ def NVVM_CpAsyncBulkTensorReduceOp :
28462845
let hasVerifier = 1;
28472846

28482847
string llvmBuilder = [{
2849-
auto [id, args, _] = NVVM::CpAsyncBulkTensorReduceOp::getIIDAndArgsWithTypes(
2848+
auto [id, args, types] = NVVM::CpAsyncBulkTensorReduceOp::getIIDAndArgsWithTypes(
28502849
*op, moduleTranslation, builder);
28512850
createIntrinsicCall(builder, id, args);
28522851
}];
@@ -2893,7 +2892,7 @@ def NVVM_CpAsyncBulkGlobalToSharedClusterOp :
28932892
}];
28942893

28952894
string llvmBuilder = [{
2896-
auto [id, args, _] = NVVM::CpAsyncBulkGlobalToSharedClusterOp::getIIDAndArgsWithTypes(
2895+
auto [id, args, types] = NVVM::CpAsyncBulkGlobalToSharedClusterOp::getIIDAndArgsWithTypes(
28972896
*op, moduleTranslation, builder);
28982897
createIntrinsicCall(builder, id, args);
28992898
}];
@@ -2982,7 +2981,7 @@ def NVVM_CpAsyncBulkSharedCTAToGlobalOp :
29822981
llvm::IRBuilderBase& builder);
29832982
}];
29842983
string llvmBuilder = [{
2985-
auto [id, args, _] = NVVM::CpAsyncBulkSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2984+
auto [id, args, types] = NVVM::CpAsyncBulkSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
29862985
*op, moduleTranslation, builder);
29872986
createIntrinsicCall(builder, id, args);
29882987
}];
@@ -3288,7 +3287,7 @@ def NVVM_MatchSyncOp : NVVM_Op<"match.sync">,
32883287
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
32893288
}];
32903289
string llvmBuilder = [{
3291-
auto [intId, args, _] =
3290+
auto [intId, args, types] =
32923291
NVVM::MatchSyncOp::getIIDAndArgsWithTypes(
32933292
*op, moduleTranslation, builder);
32943293
$res = createIntrinsicCall(builder, intId, args);
@@ -3321,7 +3320,7 @@ def NVVM_BulkStoreOp: NVVM_Op<"st.bulk"> {
33213320
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
33223321
}];
33233322
string llvmBuilder = [{
3324-
auto [intId, args, _] =
3323+
auto [intId, args, types] =
33253324
NVVM::BulkStoreOp::getIIDAndArgsWithTypes(
33263325
*op, moduleTranslation, builder);
33273326
createIntrinsicCall(builder, intId, args);
@@ -3413,7 +3412,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMa<[100, 101]>]
34133412
llvm::IRBuilderBase &builder);
34143413
}];
34153414
string llvmBuilder = [{
3416-
auto [id, args, _] = NVVM::Tcgen05AllocOp::getIIDAndArgsWithTypes(
3415+
auto [id, args, types] = NVVM::Tcgen05AllocOp::getIIDAndArgsWithTypes(
34173416
*op, moduleTranslation, builder);
34183417
createIntrinsicCall(builder, id, args);
34193418
}];
@@ -3440,7 +3439,7 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMa<[100, 10
34403439
llvm::IRBuilderBase &builder);
34413440
}];
34423441
string llvmBuilder = [{
3443-
auto [id, args, _] = NVVM::Tcgen05DeallocOp::getIIDAndArgsWithTypes(
3442+
auto [id, args, types] = NVVM::Tcgen05DeallocOp::getIIDAndArgsWithTypes(
34443443
*op, moduleTranslation, builder);
34453444
createIntrinsicCall(builder, id, args);
34463445
}];
@@ -3544,7 +3543,7 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMa<[100, 101]
35443543
}];
35453544

35463545
string llvmBuilder = [{
3547-
auto [id, args, _] = NVVM::Tcgen05CommitOp::getIIDAndArgsWithTypes(
3546+
auto [id, args, types] = NVVM::Tcgen05CommitOp::getIIDAndArgsWithTypes(
35483547
*op, moduleTranslation, builder);
35493548
createIntrinsicCall(builder, id, args);
35503549
}];
@@ -3654,7 +3653,7 @@ def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMa<[100, 101]>]> {
36543653
}];
36553654

36563655
string llvmBuilder = [{
3657-
auto [id, args, _] = NVVM::Tcgen05CpOp::getIIDAndArgsWithTypes(*op,
3656+
auto [id, args, types] = NVVM::Tcgen05CpOp::getIIDAndArgsWithTypes(*op,
36583657
moduleTranslation, builder);
36593658
createIntrinsicCall(builder, id, args);
36603659
}];
@@ -3827,7 +3826,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
38273826
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
38283827
}];
38293828
string llvmBuilder = [{
3830-
auto [id, args, _] =
3829+
auto [id, args, types] =
38313830
NVVM::Tcgen05LdOp::getIIDAndArgsWithTypes(*op,
38323831
moduleTranslation, builder);
38333832
$res = createIntrinsicCall(builder, id, args);
@@ -3907,7 +3906,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
39073906
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
39083907
}];
39093908
string llvmBuilder = [{
3910-
auto [id, args, _] =
3909+
auto [id, args, types] =
39113910
NVVM::Tcgen05StOp::getIIDAndArgsWithTypes(*op,
39123911
moduleTranslation, builder);
39133912
createIntrinsicCall(builder, id, args);
@@ -3974,7 +3973,7 @@ def NVVM_DotAccumulate4WayOp : NVVM_Op<"dot.accumulate.4way"> {
39743973
}];
39753974

39763975
string llvmBuilder = [{
3977-
auto [id, args, _] = NVVM::DotAccumulate4WayOp::getIIDAndArgsWithTypes(
3976+
auto [id, args, types] = NVVM::DotAccumulate4WayOp::getIIDAndArgsWithTypes(
39783977
*op, moduleTranslation, builder);
39793978
$res = createIntrinsicCall(builder, id, args);
39803979
}];
@@ -4028,7 +4027,7 @@ def NVVM_DotAccumulate2WayOp : NVVM_Op<"dot.accumulate.2way"> {
40284027
}];
40294028

40304029
string llvmBuilder = [{
4031-
auto [id, args, _] = NVVM::DotAccumulate2WayOp::getIIDAndArgsWithTypes(
4030+
auto [id, args, types] = NVVM::DotAccumulate2WayOp::getIIDAndArgsWithTypes(
40324031
*op, moduleTranslation, builder);
40334032
$res = createIntrinsicCall(builder, id, args);
40344033
}];

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

Lines changed: 12 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2351,7 +2351,7 @@ LdMatrixOp::getIIDAndArgsWithTypes(Operation &op, LLVM::ModuleTranslation &mt,
23512351

23522352
llvm::SmallVector<llvm::Type *> types = {args[0]->getType()};
23532353

2354-
llvm::Intrinsic::ID id;
2354+
llvm::Intrinsic::ID id = llvm::Intrinsic::not_intrinsic;
23552355

23562356
NVVM::MMALayout layout = thisOp.getLayout();
23572357
int32_t num = thisOp.getNum();
@@ -2441,6 +2441,8 @@ LdMatrixOp::getIIDAndArgsWithTypes(Operation &op, LLVM::ModuleTranslation &mt,
24412441
break;
24422442
}
24432443
}
2444+
} else {
2445+
llvm_unreachable("unknown ldmatrix kind");
24442446
}
24452447

24462448
return {id, std::move(args), types};
@@ -2458,7 +2460,7 @@ StMatrixOp::getIIDAndArgsWithTypes(Operation &op, LLVM::ModuleTranslation &mt,
24582460

24592461
llvm::SmallVector<llvm::Type *> types = {args[0]->getType()};
24602462

2461-
llvm::Intrinsic::ID id;
2463+
llvm::Intrinsic::ID id = llvm::Intrinsic::not_intrinsic;
24622464

24632465
NVVM::MMALayout layout = thisOp.getLayout();
24642466
int32_t num = thisOp.getSources().size();
@@ -2494,6 +2496,8 @@ StMatrixOp::getIIDAndArgsWithTypes(Operation &op, LLVM::ModuleTranslation &mt,
24942496
id = llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m16n8_x4_trans_b8;
24952497
break;
24962498
}
2499+
} else {
2500+
llvm_unreachable("unknown stmatrix kind");
24972501
}
24982502

24992503
return {id, std::move(args), types};
@@ -2528,7 +2532,7 @@ NVVM::IIDArgsWithTypes FenceProxyAcquireOp::getIIDAndArgsWithTypes(
25282532
args.push_back(mt.lookupValue(thisOp.getAddr()));
25292533
args.push_back(mt.lookupValue(thisOp.getSize()));
25302534

2531-
llvm::Intrinsic::ID id;
2535+
llvm::Intrinsic::ID id = llvm::Intrinsic::not_intrinsic;
25322536

25332537
NVVM::ProxyKind fromProxy = thisOp.getFromProxy();
25342538
NVVM::ProxyKind toProxy = thisOp.getToProxy();
@@ -2550,6 +2554,8 @@ NVVM::IIDArgsWithTypes FenceProxyAcquireOp::getIIDAndArgsWithTypes(
25502554
id = llvm::Intrinsic::nvvm_fence_proxy_tensormap_generic_acquire_sys;
25512555
break;
25522556
}
2557+
} else {
2558+
llvm_unreachable("unsupported proxy kinds");
25532559
}
25542560

25552561
return {id, std::move(args), {}};
@@ -2559,7 +2565,7 @@ NVVM::IIDArgsWithTypes FenceProxyReleaseOp::getIIDAndArgsWithTypes(
25592565
Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
25602566
auto thisOp = cast<NVVM::FenceProxyReleaseOp>(op);
25612567

2562-
llvm::Intrinsic::ID id;
2568+
llvm::Intrinsic::ID id = llvm::Intrinsic::not_intrinsic;
25632569

25642570
NVVM::ProxyKind fromProxy = thisOp.getFromProxy();
25652571
NVVM::ProxyKind toProxy = thisOp.getToProxy();
@@ -2581,6 +2587,8 @@ NVVM::IIDArgsWithTypes FenceProxyReleaseOp::getIIDAndArgsWithTypes(
25812587
id = llvm::Intrinsic::nvvm_fence_proxy_tensormap_generic_release_sys;
25822588
break;
25832589
}
2590+
} else {
2591+
llvm_unreachable("unsupported proxy kinds");
25842592
}
25852593

25862594
return {id, {}, {}};

0 commit comments

Comments
 (0)