@@ -416,7 +416,7 @@ def NVVM_ReduxOp :
416
416
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
417
417
}];
418
418
string llvmBuilder = [{
419
- auto [id, args, types ] =
419
+ auto [id, args, _ ] =
420
420
NVVM::ReduxOp::getIIDAndArgsWithTypes(
421
421
*op, moduleTranslation, builder);
422
422
$res = createIntrinsicCall(builder, id, args);
@@ -890,7 +890,7 @@ def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
890
890
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
891
891
}];
892
892
let llvmBuilder = [{
893
- auto [intId, args, types ] =
893
+ auto [intId, args, _ ] =
894
894
NVVM::FenceProxyAcquireOp::getIIDAndArgsWithTypes(
895
895
*op, moduleTranslation, builder);
896
896
createIntrinsicCall(builder, intId, args);
@@ -924,9 +924,10 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
924
924
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
925
925
}];
926
926
let llvmBuilder = [{
927
- auto [intId, args, types] = NVVM::FenceProxyReleaseOp::getIIDAndArgsWithTypes(
927
+ auto [intId, _, _] =
928
+ NVVM::FenceProxyReleaseOp::getIIDAndArgsWithTypes(
928
929
*op, moduleTranslation, builder);
929
- createIntrinsicCall(builder, intId, args );
930
+ createIntrinsicCall(builder, intId);
930
931
}];
931
932
932
933
let hasVerifier = 1;
@@ -1011,7 +1012,7 @@ def NVVM_ShflOp :
1011
1012
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
1012
1013
}];
1013
1014
string llvmBuilder = [{
1014
- auto [intId, args, types ] = NVVM::ShflOp::getIIDAndArgsWithTypes(
1015
+ auto [intId, args, _ ] = NVVM::ShflOp::getIIDAndArgsWithTypes(
1015
1016
*op, moduleTranslation, builder);
1016
1017
$res = createIntrinsicCall(builder, intId, args);
1017
1018
}];
@@ -1065,7 +1066,7 @@ def NVVM_VoteSyncOp
1065
1066
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
1066
1067
}];
1067
1068
string llvmBuilder = [{
1068
- auto [intId, args, types ] =
1069
+ auto [intId, args, _ ] =
1069
1070
NVVM::VoteSyncOp::getIIDAndArgsWithTypes(
1070
1071
*op, moduleTranslation, builder);
1071
1072
$res = createIntrinsicCall(builder, intId, args);
@@ -1144,7 +1145,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
1144
1145
llvm::IRBuilderBase &builder);
1145
1146
}];
1146
1147
string llvmBuilder = [{
1147
- auto [id, args, types ] = NVVM::CpAsyncOp::getIIDAndArgsWithTypes(
1148
+ auto [id, args, _ ] = NVVM::CpAsyncOp::getIIDAndArgsWithTypes(
1148
1149
*op, moduleTranslation, builder);
1149
1150
createIntrinsicCall(builder, id, args);
1150
1151
}];
@@ -2607,7 +2608,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
2607
2608
let hasVerifier = 1;
2608
2609
2609
2610
string llvmBuilder = [{
2610
- auto [id, args, types ] = NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2611
+ auto [id, args, _ ] = NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2611
2612
*op, moduleTranslation, builder);
2612
2613
createIntrinsicCall(builder, id, args);
2613
2614
}];
@@ -2685,7 +2686,7 @@ def NVVM_PrefetchOp : NVVM_Op<"prefetch",
2685
2686
}
2686
2687
}];
2687
2688
let llvmBuilder = [{
2688
- auto [id, args, types ] = NVVM::PrefetchOp::getIIDAndArgsWithTypes(*op,
2689
+ auto [id, args, _ ] = NVVM::PrefetchOp::getIIDAndArgsWithTypes(*op,
2689
2690
moduleTranslation, builder);
2690
2691
2691
2692
if(op.getTensormap())
@@ -2733,7 +2734,7 @@ def NVVM_CpAsyncBulkPrefetchOp : NVVM_Op<"cp.async.bulk.prefetch"> {
2733
2734
}];
2734
2735
2735
2736
string llvmBuilder = [{
2736
- auto [id, args, types ] = NVVM::CpAsyncBulkPrefetchOp::getIIDAndArgsWithTypes(
2737
+ auto [id, args, _ ] = NVVM::CpAsyncBulkPrefetchOp::getIIDAndArgsWithTypes(
2737
2738
*op, moduleTranslation, builder);
2738
2739
createIntrinsicCall(builder, id, args);
2739
2740
}];
@@ -2776,7 +2777,7 @@ def NVVM_CpAsyncBulkTensorPrefetchOp :
2776
2777
let hasVerifier = 1;
2777
2778
2778
2779
string llvmBuilder = [{
2779
- auto [id, args, types ] = NVVM::CpAsyncBulkTensorPrefetchOp::getIIDAndArgsWithTypes(
2780
+ auto [id, args, _ ] = NVVM::CpAsyncBulkTensorPrefetchOp::getIIDAndArgsWithTypes(
2780
2781
*op, moduleTranslation, builder);
2781
2782
createIntrinsicCall(builder, id, args);
2782
2783
}];
@@ -2845,7 +2846,7 @@ def NVVM_CpAsyncBulkTensorReduceOp :
2845
2846
let hasVerifier = 1;
2846
2847
2847
2848
string llvmBuilder = [{
2848
- auto [id, args, types ] = NVVM::CpAsyncBulkTensorReduceOp::getIIDAndArgsWithTypes(
2849
+ auto [id, args, _ ] = NVVM::CpAsyncBulkTensorReduceOp::getIIDAndArgsWithTypes(
2849
2850
*op, moduleTranslation, builder);
2850
2851
createIntrinsicCall(builder, id, args);
2851
2852
}];
@@ -2892,7 +2893,7 @@ def NVVM_CpAsyncBulkGlobalToSharedClusterOp :
2892
2893
}];
2893
2894
2894
2895
string llvmBuilder = [{
2895
- auto [id, args, types ] = NVVM::CpAsyncBulkGlobalToSharedClusterOp::getIIDAndArgsWithTypes(
2896
+ auto [id, args, _ ] = NVVM::CpAsyncBulkGlobalToSharedClusterOp::getIIDAndArgsWithTypes(
2896
2897
*op, moduleTranslation, builder);
2897
2898
createIntrinsicCall(builder, id, args);
2898
2899
}];
@@ -2981,7 +2982,7 @@ def NVVM_CpAsyncBulkSharedCTAToGlobalOp :
2981
2982
llvm::IRBuilderBase& builder);
2982
2983
}];
2983
2984
string llvmBuilder = [{
2984
- auto [id, args, types ] = NVVM::CpAsyncBulkSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2985
+ auto [id, args, _ ] = NVVM::CpAsyncBulkSharedCTAToGlobalOp::getIIDAndArgsWithTypes(
2985
2986
*op, moduleTranslation, builder);
2986
2987
createIntrinsicCall(builder, id, args);
2987
2988
}];
@@ -3287,7 +3288,7 @@ def NVVM_MatchSyncOp : NVVM_Op<"match.sync">,
3287
3288
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
3288
3289
}];
3289
3290
string llvmBuilder = [{
3290
- auto [intId, args, types ] =
3291
+ auto [intId, args, _ ] =
3291
3292
NVVM::MatchSyncOp::getIIDAndArgsWithTypes(
3292
3293
*op, moduleTranslation, builder);
3293
3294
$res = createIntrinsicCall(builder, intId, args);
@@ -3320,7 +3321,7 @@ def NVVM_BulkStoreOp: NVVM_Op<"st.bulk"> {
3320
3321
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
3321
3322
}];
3322
3323
string llvmBuilder = [{
3323
- auto [intId, args, types ] =
3324
+ auto [intId, args, _ ] =
3324
3325
NVVM::BulkStoreOp::getIIDAndArgsWithTypes(
3325
3326
*op, moduleTranslation, builder);
3326
3327
createIntrinsicCall(builder, intId, args);
@@ -3412,7 +3413,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMa<[100, 101]>]
3412
3413
llvm::IRBuilderBase &builder);
3413
3414
}];
3414
3415
string llvmBuilder = [{
3415
- auto [id, args, types ] = NVVM::Tcgen05AllocOp::getIIDAndArgsWithTypes(
3416
+ auto [id, args, _ ] = NVVM::Tcgen05AllocOp::getIIDAndArgsWithTypes(
3416
3417
*op, moduleTranslation, builder);
3417
3418
createIntrinsicCall(builder, id, args);
3418
3419
}];
@@ -3439,7 +3440,7 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMa<[100, 10
3439
3440
llvm::IRBuilderBase &builder);
3440
3441
}];
3441
3442
string llvmBuilder = [{
3442
- auto [id, args, types ] = NVVM::Tcgen05DeallocOp::getIIDAndArgsWithTypes(
3443
+ auto [id, args, _ ] = NVVM::Tcgen05DeallocOp::getIIDAndArgsWithTypes(
3443
3444
*op, moduleTranslation, builder);
3444
3445
createIntrinsicCall(builder, id, args);
3445
3446
}];
@@ -3543,7 +3544,7 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMa<[100, 101]
3543
3544
}];
3544
3545
3545
3546
string llvmBuilder = [{
3546
- auto [id, args, types ] = NVVM::Tcgen05CommitOp::getIIDAndArgsWithTypes(
3547
+ auto [id, args, _ ] = NVVM::Tcgen05CommitOp::getIIDAndArgsWithTypes(
3547
3548
*op, moduleTranslation, builder);
3548
3549
createIntrinsicCall(builder, id, args);
3549
3550
}];
@@ -3653,7 +3654,7 @@ def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMa<[100, 101]>]> {
3653
3654
}];
3654
3655
3655
3656
string llvmBuilder = [{
3656
- auto [id, args, types ] = NVVM::Tcgen05CpOp::getIIDAndArgsWithTypes(*op,
3657
+ auto [id, args, _ ] = NVVM::Tcgen05CpOp::getIIDAndArgsWithTypes(*op,
3657
3658
moduleTranslation, builder);
3658
3659
createIntrinsicCall(builder, id, args);
3659
3660
}];
@@ -3826,7 +3827,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
3826
3827
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
3827
3828
}];
3828
3829
string llvmBuilder = [{
3829
- auto [id, args, types ] =
3830
+ auto [id, args, _ ] =
3830
3831
NVVM::Tcgen05LdOp::getIIDAndArgsWithTypes(*op,
3831
3832
moduleTranslation, builder);
3832
3833
$res = createIntrinsicCall(builder, id, args);
@@ -3906,7 +3907,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
3906
3907
LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
3907
3908
}];
3908
3909
string llvmBuilder = [{
3909
- auto [id, args, types ] =
3910
+ auto [id, args, _ ] =
3910
3911
NVVM::Tcgen05StOp::getIIDAndArgsWithTypes(*op,
3911
3912
moduleTranslation, builder);
3912
3913
createIntrinsicCall(builder, id, args);
@@ -3973,7 +3974,7 @@ def NVVM_DotAccumulate4WayOp : NVVM_Op<"dot.accumulate.4way"> {
3973
3974
}];
3974
3975
3975
3976
string llvmBuilder = [{
3976
- auto [id, args, types ] = NVVM::DotAccumulate4WayOp::getIIDAndArgsWithTypes(
3977
+ auto [id, args, _ ] = NVVM::DotAccumulate4WayOp::getIIDAndArgsWithTypes(
3977
3978
*op, moduleTranslation, builder);
3978
3979
$res = createIntrinsicCall(builder, id, args);
3979
3980
}];
@@ -4027,7 +4028,7 @@ def NVVM_DotAccumulate2WayOp : NVVM_Op<"dot.accumulate.2way"> {
4027
4028
}];
4028
4029
4029
4030
string llvmBuilder = [{
4030
- auto [id, args, types ] = NVVM::DotAccumulate2WayOp::getIIDAndArgsWithTypes(
4031
+ auto [id, args, _ ] = NVVM::DotAccumulate2WayOp::getIIDAndArgsWithTypes(
4031
4032
*op, moduleTranslation, builder);
4032
4033
$res = createIntrinsicCall(builder, id, args);
4033
4034
}];
0 commit comments