Skip to content

Commit c7ee204

Browse files
sushgokhKugan Vivekanandarajah
sushgokh
and
Kugan Vivekanandarajah
authored
[OpenMP] Fix stack corruption due to argument mismatch (#96386)
While lowering (#pragma omp target update from), clang's generated .omp_task_entry. is setting up 9 arguments while calling __tgt_target_data_update_nowait_mapper. At the same time, in __tgt_target_data_update_nowait_mapper, call to targetData<TaskAsyncInfoWrapperTy>() is converted to a sibcall assuming it has the argument count listed in the signature. AARCH64 asm sequence for this is as follows (removed unrelated insns): ` .omp_task_entry..108: sub sp, sp, #32 stp x29, x30, sp, #16 // 16-byte Folded Spill add x29, sp, #16 str x8, sp, #8. // stack canary str xzr, [sp] bl __tgt_target_data_update_nowait_mapper __tgt_target_data_update_nowait_mapper: sub sp, sp, #32 stp x29, x30, sp, #16 // 16-byte Folded Spill add x29, sp, #16 str x8, sp, #8 // stack canary // Sibcall argument setup adrp x8, :got:_Z16targetDataUpdateP7ident_tR8DeviceTyiPPvS4_PlS5_S4_S4_R11AsyncInfoTyb ldr x8, [x8, :got_lo12:_Z16targetDataUpdateP7ident_tR8DeviceTyiPPvS4_PlS5_S4_S4_R11AsyncInfoTyb] stp x9, x8, x29, #16 adrp x8, .L.str.8 add x8, x8, :lo12:.L.str.8 str x8, x29, #32. <==. This is the insn that erases $fp ldp x29, x30, sp, #16 // 16-byte Folded Reload add sp, sp, #32 // Sibcall b ZL10targetDataI22TaskAsyncInfoWrapperTyEvP7ident_tliPPvS4_PlS5_S4_S4_PFiS2_R8DeviceTyiS4_S4_S5_S5_S4_S4_R11AsyncInfoTybEPKcSD ` On AArch64, call to __tgt_target_data_update_nowait_mapper in .omp_task_entry. sets up only single space on stack and this results in ovewriting $fp and subsequent stack corruption. This issue can be credited to discrepancy of __tgt_target_data_update_nowait_mapper signature in openmp/libomptarget/include/omptarget.h taking 13 arguments while clang/lib/CodeGen/CGOpenMPRuntime.cpp and llvm/include/llvm/Frontend/OpenMP/OMPKinds.def taking only 9 arguments. This patch modifies __tgt_target_data_update_nowait_mapper signature to match .omp_task_entry usage(and other 2 files mentioned above). Co-authored-by: Kugan Vivekanandarajah <kvivekananda@nvidia.com>
1 parent db782b4 commit c7ee204

File tree

7 files changed

+50
-39
lines changed

7 files changed

+50
-39
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

+13-10
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,7 @@
3030
#include "llvm/ADT/ArrayRef.h"
3131
#include "llvm/ADT/SetOperations.h"
3232
#include "llvm/ADT/SmallBitVector.h"
33+
#include "llvm/ADT/SmallVector.h"
3334
#include "llvm/ADT/StringExtras.h"
3435
#include "llvm/Bitcode/BitcodeReader.h"
3536
#include "llvm/IR/Constants.h"
@@ -10357,16 +10358,12 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
1035710358
// Source location for the ident struct
1035810359
llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
1035910360

10360-
llvm::Value *OffloadingArgs[] = {
10361-
RTLoc,
10362-
DeviceID,
10363-
PointerNum,
10364-
InputInfo.BasePointersArray.emitRawPointer(CGF),
10365-
InputInfo.PointersArray.emitRawPointer(CGF),
10366-
InputInfo.SizesArray.emitRawPointer(CGF),
10367-
MapTypesArray,
10368-
MapNamesArray,
10369-
InputInfo.MappersArray.emitRawPointer(CGF)};
10361+
SmallVector<llvm::Value *, 13> OffloadingArgs(
10362+
{RTLoc, DeviceID, PointerNum,
10363+
InputInfo.BasePointersArray.emitRawPointer(CGF),
10364+
InputInfo.PointersArray.emitRawPointer(CGF),
10365+
InputInfo.SizesArray.emitRawPointer(CGF), MapTypesArray, MapNamesArray,
10366+
InputInfo.MappersArray.emitRawPointer(CGF)});
1037010367

1037110368
// Select the right runtime function call for each standalone
1037210369
// directive.
@@ -10455,6 +10452,12 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
1045510452
llvm_unreachable("Unexpected standalone target data directive.");
1045610453
break;
1045710454
}
10455+
if (HasNowait) {
10456+
OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.Int32Ty));
10457+
OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.VoidPtrTy));
10458+
OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.Int32Ty));
10459+
OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.VoidPtrTy));
10460+
}
1045810461
CGF.EmitRuntimeCall(
1045910462
OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), RTLFn),
1046010463
OffloadingArgs);

‎clang/test/OpenMP/declare_mapper_codegen.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -514,7 +514,7 @@ void foo(int a){
514514
// CK0: }
515515

516516
// CK0: define internal void [[OMP_OUTLINED_16:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
517-
// CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EDNWTYPES]], ptr null, ptr [[MPR:%.+]])
517+
// CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EDNWTYPES]], ptr null, ptr [[MPR:%.+]], i32 0, ptr null, i32 0, ptr null)
518518
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
519519
// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
520520
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
@@ -533,7 +533,7 @@ void foo(int a){
533533
// CK0: }
534534

535535
// CK0: define internal void [[OMP_OUTLINED_23:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
536-
// CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EXDNWTYPES]], ptr null, ptr [[MPR:%.+]])
536+
// CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EXDNWTYPES]], ptr null, ptr [[MPR:%.+]], i32 0, ptr null, i32 0, ptr null)
537537
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
538538
// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
539539
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
@@ -551,7 +551,7 @@ void foo(int a){
551551
// CK0: }
552552

553553
// CK0: define internal void [[OMP_OUTLINED_32:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
554-
// CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[FNWTYPES]], ptr null, ptr [[MPR:%.+]])
554+
// CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[FNWTYPES]], ptr null, ptr [[MPR:%.+]], i32 0, ptr null, i32 0, ptr null)
555555
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
556556
// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
557557
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0

‎clang/test/OpenMP/target_enter_data_codegen.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ void foo(int arg) {
209209

210210

211211
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
212-
// CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BPADDR:%[^,]+]], ptr [[PADDR:%[^,]+]], ptr [[SZADDR:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
212+
// CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BPADDR:%[^,]+]], ptr [[PADDR:%[^,]+]], ptr [[SZADDR:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null, i32 0, ptr null, i32 0, ptr null)
213213
// CK1-DAG: [[BPADDR]] = load ptr, ptr [[FPBP:%[^,]+]], align
214214
// CK1-DAG: [[PADDR]] = load ptr, ptr [[FPP:%[^,]+]], align
215215
// CK1-DAG: [[SZADDR]] = load ptr, ptr [[FPSZ:%[^,]+]], align

‎clang/test/OpenMP/target_exit_data_codegen.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -206,7 +206,7 @@ void foo(int arg) {
206206
}
207207

208208
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}})
209-
// CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
209+
// CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null, i32 0, ptr null, i32 0, ptr null)
210210
// CK1-DAG: [[BP]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align
211211
// CK1-DAG: [[P]] = load ptr, ptr [[FPPADDR:%[^,]+]], align
212212
// CK1-DAG: [[SZ]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align

‎clang/test/OpenMP/target_update_codegen.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -151,7 +151,7 @@ void foo(int arg) {
151151
}
152152

153153
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}})
154-
// CK1-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
154+
// CK1-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null, i32 0, ptr null, i32 0, ptr null)
155155
// CK1-DAG: [[BP]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align
156156
// CK1-DAG: [[P]] = load ptr, ptr [[FPPADDR:%[^,]+]], align
157157
// CK1-DAG: [[SZ]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align

‎llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

+19-11
Original file line numberDiff line numberDiff line change
@@ -439,19 +439,22 @@ __OMP_RTL(__tgt_target_kernel_nowait, false, Int32, IdentPtr, Int64, Int32,
439439
Int32, VoidPtr, KernelArgsPtr, Int32, VoidPtr, Int32, VoidPtr)
440440
__OMP_RTL(__tgt_target_data_begin_mapper, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr,
441441
VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
442-
__OMP_RTL(__tgt_target_data_begin_nowait_mapper, false, Void, IdentPtr, Int64, Int32,
443-
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
442+
__OMP_RTL(__tgt_target_data_begin_nowait_mapper, false, Void, IdentPtr, Int64,
443+
Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr,
444+
VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr)
444445
__OMP_RTL(__tgt_target_data_begin_mapper_issue, false, Void, IdentPtr, Int64, Int32,
445446
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr, AsyncInfoPtr)
446447
__OMP_RTL(__tgt_target_data_begin_mapper_wait, false, Void, Int64, AsyncInfoPtr)
447448
__OMP_RTL(__tgt_target_data_end_mapper, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr,
448449
VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
449-
__OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, IdentPtr, Int64, Int32,
450-
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
450+
__OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, IdentPtr, Int64,
451+
Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr,
452+
VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr)
451453
__OMP_RTL(__tgt_target_data_update_mapper, false, Void, IdentPtr, Int64, Int32,
452454
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
453-
__OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, IdentPtr, Int64, Int32,
454-
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
455+
__OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, IdentPtr, Int64,
456+
Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr,
457+
VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr)
455458
__OMP_RTL(__tgt_mapper_num_components, false, Int64, VoidPtr)
456459
__OMP_RTL(__tgt_push_mapper_component, false, Void, VoidPtr, VoidPtr, VoidPtr,
457460
Int64, Int64, VoidPtr)
@@ -1029,10 +1032,12 @@ __OMP_RTL_ATTRS(__tgt_target_kernel_nowait, ForkAttrs, SExt,
10291032
SExt))
10301033
__OMP_RTL_ATTRS(__tgt_target_data_begin_mapper, ForkAttrs, AttributeSet(),
10311034
ParamAttrs(AttributeSet(), AttributeSet(), SExt))
1032-
__OMP_RTL_ATTRS(__tgt_target_data_begin_nowait_mapper, ForkAttrs, AttributeSet(),
1035+
__OMP_RTL_ATTRS(__tgt_target_data_begin_nowait_mapper, ForkAttrs,
1036+
AttributeSet(),
10331037
ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(),
10341038
AttributeSet(), AttributeSet(), AttributeSet(),
1035-
AttributeSet(), AttributeSet()))
1039+
AttributeSet(), AttributeSet(), SExt, AttributeSet(),
1040+
SExt, AttributeSet()))
10361041
__OMP_RTL_ATTRS(__tgt_target_data_begin_mapper_issue, AttributeSet(),
10371042
AttributeSet(),
10381043
ParamAttrs(AttributeSet(), AttributeSet(), SExt))
@@ -1041,13 +1046,16 @@ __OMP_RTL_ATTRS(__tgt_target_data_end_mapper, ForkAttrs, AttributeSet(),
10411046
__OMP_RTL_ATTRS(__tgt_target_data_end_nowait_mapper, ForkAttrs, AttributeSet(),
10421047
ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(),
10431048
AttributeSet(), AttributeSet(), AttributeSet(),
1044-
AttributeSet(), AttributeSet()))
1049+
AttributeSet(), AttributeSet(), SExt, AttributeSet(),
1050+
SExt, AttributeSet()))
10451051
__OMP_RTL_ATTRS(__tgt_target_data_update_mapper, ForkAttrs, AttributeSet(),
10461052
ParamAttrs(AttributeSet(), AttributeSet(), SExt))
1047-
__OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs, AttributeSet(),
1053+
__OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs,
1054+
AttributeSet(),
10481055
ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(),
10491056
AttributeSet(), AttributeSet(), AttributeSet(),
1050-
AttributeSet(), AttributeSet()))
1057+
AttributeSet(), AttributeSet(), SExt, AttributeSet(),
1058+
SExt, AttributeSet()))
10511059
__OMP_RTL_ATTRS(__tgt_mapper_num_components, ForkAttrs, AttributeSet(),
10521060
ParamAttrs())
10531061
__OMP_RTL_ATTRS(__tgt_push_mapper_component, ForkAttrs, AttributeSet(),

‎llvm/test/Transforms/OpenMP/add_attributes.ll

+12-12
Original file line numberDiff line numberDiff line change
@@ -643,15 +643,15 @@ declare i32 @__tgt_target_teams_nowait_mapper(ptr, i64, ptr, i32, ptr, ptr, ptr,
643643

644644
declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
645645

646-
declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
646+
declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
647647

648648
declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
649649

650-
declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
650+
declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
651651

652652
declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
653653

654-
declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
654+
declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
655655

656656
declare i64 @__tgt_mapper_num_components(ptr)
657657

@@ -1250,19 +1250,19 @@ declare i32 @__tgt_target_kernel_nowait(ptr, i64, i32, i32, ptr, ptr, i32, ptr,
12501250
; CHECK-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
12511251

12521252
; CHECK: ; Function Attrs: nounwind
1253-
; CHECK-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1253+
; CHECK-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
12541254

12551255
; CHECK: ; Function Attrs: nounwind
12561256
; CHECK-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
12571257

12581258
; CHECK: ; Function Attrs: nounwind
1259-
; CHECK-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1259+
; CHECK-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
12601260

12611261
; CHECK: ; Function Attrs: nounwind
12621262
; CHECK-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
12631263

12641264
; CHECK: ; Function Attrs: nounwind
1265-
; CHECK-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1265+
; CHECK-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
12661266

12671267
; CHECK: ; Function Attrs: nounwind
12681268
; CHECK-NEXT: declare i64 @__tgt_mapper_num_components(ptr)
@@ -1892,19 +1892,19 @@ declare i32 @__tgt_target_kernel_nowait(ptr, i64, i32, i32, ptr, ptr, i32, ptr,
18921892
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
18931893

18941894
; OPTIMISTIC: ; Function Attrs: nounwind
1895-
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1895+
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
18961896

18971897
; OPTIMISTIC: ; Function Attrs: nounwind
18981898
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
18991899

19001900
; OPTIMISTIC: ; Function Attrs: nounwind
1901-
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1901+
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
19021902

19031903
; OPTIMISTIC: ; Function Attrs: nounwind
19041904
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
19051905

19061906
; OPTIMISTIC: ; Function Attrs: nounwind
1907-
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
1907+
; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr)
19081908

19091909
; OPTIMISTIC: ; Function Attrs: nounwind
19101910
; OPTIMISTIC-NEXT: declare i64 @__tgt_mapper_num_components(ptr)
@@ -2547,19 +2547,19 @@ declare i32 @__tgt_target_kernel_nowait(ptr, i64, i32, i32, ptr, ptr, i32, ptr,
25472547
; EXT-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
25482548

25492549
; EXT: ; Function Attrs: nounwind
2550-
; EXT-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
2550+
; EXT-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr, i32 signext, ptr, i32 signext, ptr)
25512551

25522552
; EXT: ; Function Attrs: nounwind
25532553
; EXT-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
25542554

25552555
; EXT: ; Function Attrs: nounwind
2556-
; EXT-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
2556+
; EXT-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr, i32 signext, ptr, i32 signext, ptr)
25572557

25582558
; EXT: ; Function Attrs: nounwind
25592559
; EXT-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
25602560

25612561
; EXT: ; Function Attrs: nounwind
2562-
; EXT-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr)
2562+
; EXT-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr, i32 signext, ptr, i32 signext, ptr)
25632563

25642564
; EXT: ; Function Attrs: nounwind
25652565
; EXT-NEXT: declare i64 @__tgt_mapper_num_components(ptr)

0 commit comments

Comments
 (0)