From c7ee20433c43e45658031a340e221097a32a469f Mon Sep 17 00:00:00 2001 From: Sushant Gokhale Date: Fri, 5 Jul 2024 10:39:15 +0530 Subject: [PATCH] [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() 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 --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 23 +++++++------- clang/test/OpenMP/declare_mapper_codegen.cpp | 6 ++-- .../test/OpenMP/target_enter_data_codegen.cpp | 2 +- .../test/OpenMP/target_exit_data_codegen.cpp | 2 +- clang/test/OpenMP/target_update_codegen.cpp | 2 +- .../include/llvm/Frontend/OpenMP/OMPKinds.def | 30 ++++++++++++------- llvm/test/Transforms/OpenMP/add_attributes.ll | 24 +++++++-------- 7 files changed, 50 insertions(+), 39 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 3febeed7b72d9b..8bc202f402aa39 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -30,6 +30,7 @@ #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/SetOperations.h" #include "llvm/ADT/SmallBitVector.h" +#include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Bitcode/BitcodeReader.h" #include "llvm/IR/Constants.h" @@ -10357,16 +10358,12 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall( // Source location for the ident struct llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc()); - llvm::Value *OffloadingArgs[] = { - RTLoc, - DeviceID, - PointerNum, - InputInfo.BasePointersArray.emitRawPointer(CGF), - InputInfo.PointersArray.emitRawPointer(CGF), - InputInfo.SizesArray.emitRawPointer(CGF), - MapTypesArray, - MapNamesArray, - InputInfo.MappersArray.emitRawPointer(CGF)}; + SmallVector OffloadingArgs( + {RTLoc, DeviceID, PointerNum, + InputInfo.BasePointersArray.emitRawPointer(CGF), + InputInfo.PointersArray.emitRawPointer(CGF), + InputInfo.SizesArray.emitRawPointer(CGF), MapTypesArray, MapNamesArray, + InputInfo.MappersArray.emitRawPointer(CGF)}); // Select the right runtime function call for each standalone // directive. @@ -10455,6 +10452,12 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall( llvm_unreachable("Unexpected standalone target data directive."); break; } + if (HasNowait) { + OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.Int32Ty)); + OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.VoidPtrTy)); + OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.Int32Ty)); + OffloadingArgs.push_back(llvm::Constant::getNullValue(CGF.VoidPtrTy)); + } CGF.EmitRuntimeCall( OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), RTLFn), OffloadingArgs); diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index 647e2a09074353..8c1db22fb7ff73 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -514,7 +514,7 @@ void foo(int a){ // CK0: } // CK0: define internal void [[OMP_OUTLINED_16:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}} -// 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:%.+]]) +// 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) // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 @@ -533,7 +533,7 @@ void foo(int a){ // CK0: } // CK0: define internal void [[OMP_OUTLINED_23:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}} -// 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:%.+]]) +// 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) // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 @@ -551,7 +551,7 @@ void foo(int a){ // CK0: } // CK0: define internal void [[OMP_OUTLINED_32:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}} -// 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:%.+]]) +// 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) // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0 diff --git a/clang/test/OpenMP/target_enter_data_codegen.cpp b/clang/test/OpenMP/target_enter_data_codegen.cpp index 147d372fccaaf2..dd94020b28e11e 100644 --- a/clang/test/OpenMP/target_enter_data_codegen.cpp +++ b/clang/test/OpenMP/target_enter_data_codegen.cpp @@ -209,7 +209,7 @@ void foo(int arg) { // CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1) -// 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) +// 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) // CK1-DAG: [[BPADDR]] = load ptr, ptr [[FPBP:%[^,]+]], align // CK1-DAG: [[PADDR]] = load ptr, ptr [[FPP:%[^,]+]], align // CK1-DAG: [[SZADDR]] = load ptr, ptr [[FPSZ:%[^,]+]], align diff --git a/clang/test/OpenMP/target_exit_data_codegen.cpp b/clang/test/OpenMP/target_exit_data_codegen.cpp index 96a0a5063dd6e0..01e3b4fc3364e3 100644 --- a/clang/test/OpenMP/target_exit_data_codegen.cpp +++ b/clang/test/OpenMP/target_exit_data_codegen.cpp @@ -206,7 +206,7 @@ void foo(int arg) { } // CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}}) -// 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) +// 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) // CK1-DAG: [[BP]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align // CK1-DAG: [[P]] = load ptr, ptr [[FPPADDR:%[^,]+]], align // CK1-DAG: [[SZ]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp index b577be3c1b4966..de0d4bda86088a 100644 --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -151,7 +151,7 @@ void foo(int arg) { } // CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}}) -// 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) +// 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) // CK1-DAG: [[BP]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align // CK1-DAG: [[P]] = load ptr, ptr [[FPPADDR:%[^,]+]], align // CK1-DAG: [[SZ]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def index 535a3a74d7eaab..d9e9c14af3b157 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -439,19 +439,22 @@ __OMP_RTL(__tgt_target_kernel_nowait, false, Int32, IdentPtr, Int64, Int32, Int32, VoidPtr, KernelArgsPtr, Int32, VoidPtr, Int32, VoidPtr) __OMP_RTL(__tgt_target_data_begin_mapper, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) -__OMP_RTL(__tgt_target_data_begin_nowait_mapper, false, Void, IdentPtr, Int64, Int32, - VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) +__OMP_RTL(__tgt_target_data_begin_nowait_mapper, false, Void, IdentPtr, Int64, + Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, + VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr) __OMP_RTL(__tgt_target_data_begin_mapper_issue, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr, AsyncInfoPtr) __OMP_RTL(__tgt_target_data_begin_mapper_wait, false, Void, Int64, AsyncInfoPtr) __OMP_RTL(__tgt_target_data_end_mapper, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) -__OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, IdentPtr, Int64, Int32, - VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) +__OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, IdentPtr, Int64, + Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, + VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr) __OMP_RTL(__tgt_target_data_update_mapper, false, Void, IdentPtr, Int64, Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) -__OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, IdentPtr, Int64, Int32, - VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) +__OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, IdentPtr, Int64, + Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, + VoidPtrPtr, Int32, VoidPtr, Int32, VoidPtr) __OMP_RTL(__tgt_mapper_num_components, false, Int64, VoidPtr) __OMP_RTL(__tgt_push_mapper_component, false, Void, VoidPtr, VoidPtr, VoidPtr, Int64, Int64, VoidPtr) @@ -1029,10 +1032,12 @@ __OMP_RTL_ATTRS(__tgt_target_kernel_nowait, ForkAttrs, SExt, SExt)) __OMP_RTL_ATTRS(__tgt_target_data_begin_mapper, ForkAttrs, AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt)) -__OMP_RTL_ATTRS(__tgt_target_data_begin_nowait_mapper, ForkAttrs, AttributeSet(), +__OMP_RTL_ATTRS(__tgt_target_data_begin_nowait_mapper, ForkAttrs, + AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(), AttributeSet(), AttributeSet(), AttributeSet(), - AttributeSet(), AttributeSet())) + AttributeSet(), AttributeSet(), SExt, AttributeSet(), + SExt, AttributeSet())) __OMP_RTL_ATTRS(__tgt_target_data_begin_mapper_issue, AttributeSet(), AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt)) @@ -1041,13 +1046,16 @@ __OMP_RTL_ATTRS(__tgt_target_data_end_mapper, ForkAttrs, AttributeSet(), __OMP_RTL_ATTRS(__tgt_target_data_end_nowait_mapper, ForkAttrs, AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(), AttributeSet(), AttributeSet(), AttributeSet(), - AttributeSet(), AttributeSet())) + AttributeSet(), AttributeSet(), SExt, AttributeSet(), + SExt, AttributeSet())) __OMP_RTL_ATTRS(__tgt_target_data_update_mapper, ForkAttrs, AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt)) -__OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs, AttributeSet(), +__OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs, + AttributeSet(), ParamAttrs(AttributeSet(), AttributeSet(), SExt, AttributeSet(), AttributeSet(), AttributeSet(), AttributeSet(), - AttributeSet(), AttributeSet())) + AttributeSet(), AttributeSet(), SExt, AttributeSet(), + SExt, AttributeSet())) __OMP_RTL_ATTRS(__tgt_mapper_num_components, ForkAttrs, AttributeSet(), ParamAttrs()) __OMP_RTL_ATTRS(__tgt_push_mapper_component, ForkAttrs, AttributeSet(), diff --git a/llvm/test/Transforms/OpenMP/add_attributes.ll b/llvm/test/Transforms/OpenMP/add_attributes.ll index ebcca3067f045a..fefb1900afae7e 100644 --- a/llvm/test/Transforms/OpenMP/add_attributes.ll +++ b/llvm/test/Transforms/OpenMP/add_attributes.ll @@ -643,15 +643,15 @@ declare i32 @__tgt_target_teams_nowait_mapper(ptr, i64, ptr, i32, ptr, ptr, ptr, declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) -declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) -declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) -declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) declare i64 @__tgt_mapper_num_components(ptr) @@ -1250,19 +1250,19 @@ declare i32 @__tgt_target_kernel_nowait(ptr, i64, i32, i32, ptr, ptr, i32, ptr, ; CHECK-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; CHECK: ; Function Attrs: nounwind -; CHECK-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; CHECK-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; CHECK: ; Function Attrs: nounwind ; CHECK-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; CHECK: ; Function Attrs: nounwind -; CHECK-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; CHECK-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; CHECK: ; Function Attrs: nounwind ; CHECK-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; CHECK: ; Function Attrs: nounwind -; CHECK-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; CHECK-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; CHECK: ; Function Attrs: nounwind ; 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, ; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind -; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; OPTIMISTIC-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind ; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind -; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; OPTIMISTIC-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind ; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind -; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr) +; OPTIMISTIC-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, ptr, i32, ptr) ; OPTIMISTIC: ; Function Attrs: nounwind ; 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, ; EXT-NEXT: declare void @__tgt_target_data_begin_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) ; EXT: ; Function Attrs: nounwind -; EXT-NEXT: declare void @__tgt_target_data_begin_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) +; 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) ; EXT: ; Function Attrs: nounwind ; EXT-NEXT: declare void @__tgt_target_data_end_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) ; EXT: ; Function Attrs: nounwind -; EXT-NEXT: declare void @__tgt_target_data_end_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) +; 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) ; EXT: ; Function Attrs: nounwind ; EXT-NEXT: declare void @__tgt_target_data_update_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) ; EXT: ; Function Attrs: nounwind -; EXT-NEXT: declare void @__tgt_target_data_update_nowait_mapper(ptr, i64, i32 signext, ptr, ptr, ptr, ptr, ptr, ptr) +; 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) ; EXT: ; Function Attrs: nounwind ; EXT-NEXT: declare i64 @__tgt_mapper_num_components(ptr)