[OpenMP] Fix stack corruption due to argument mismatch #96386

sushgokh · 2024-06-22T11:50:56Z

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).

llvmbot · 2024-06-22T11:51:27Z

@llvm/pr-subscribers-llvm-transforms
@llvm/pr-subscribers-flang-openmp
@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-offload

Author: Sushant Gokhale (sushgokh)

Changes

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).

Full diff: https://github.com/llvm/llvm-project/pull/96386.diff

2 Files Affected:

(modified) offload/include/omptarget.h (+1-2)
(modified) offload/src/interface.cpp (+1-2)

diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h
index 323dee41630f2..968589e866334 100644
--- a/offload/include/omptarget.h
+++ b/offload/include/omptarget.h
@@ -388,8 +388,7 @@ void __tgt_target_data_update_mapper(ident_t *Loc, int64_t DeviceId,
 void __tgt_target_data_update_nowait_mapper(
     ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
     void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
-    void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
-    void *NoAliasDepList);
+    void **ArgMappers);
 
 // Performs the same actions as data_begin in case ArgNum is non-zero
 // and initiates run of offloaded region on target platform; if ArgNum
diff --git a/offload/src/interface.cpp b/offload/src/interface.cpp
index 763b051cc6d77..27562388c2f11 100644
--- a/offload/src/interface.cpp
+++ b/offload/src/interface.cpp
@@ -207,8 +207,7 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *Loc, int64_t DeviceId,
 EXTERN void __tgt_target_data_update_nowait_mapper(
     ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
     void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
-    void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
-    void *NoAliasDepList) {
+    void **ArgMappers) {
   OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
   targetData<TaskAsyncInfoWrapperTy>(
       Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames,

jhuber6 · 2024-06-22T11:54:02Z

offload/src/interface.cpp

@@ -207,8 +207,7 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *Loc, int64_t DeviceId,
 EXTERN void __tgt_target_data_update_nowait_mapper(
    ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
    void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
-    void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
-    void *NoAliasDepList) {
+    void **ArgMappers) {


These are all unused? Seems weird that we're not using them at all.

The original design was from libomp task interface, but we never use it.

However, I think we might want to keep them, as I definitely had a patch long time ago that uses them. A better fix would be to modify the front end to emit null value correspondingly.

However, I think we might want to keep them, as I definitely had a patch long time ago that uses them. A better fix would be to modify the front end to emit null value correspondingly.

All the nowait calls, listed below, suffer from the same discrepancy

__tgt_target_data_begin_nowait_mapper
__tgt_target_data_end_nowait_mapper
__tgt_target_data_update_nowait_mapper

while their counterparts, derived here, only have 9 arguments

So, I am thinking of modifying all the nowait calls while addressing this specific issue.

@shiltian If you still think that modifying the front-end would be better here, could you tell me what changes are exactly required in the OffloadingArgs here ? I hope this is what you are suggesting.

Yes. If it is nowait, push extra null values.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

shiltian

Besides, don't you need to update clang tests?

shiltian

LGTM

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, llvm#32 stp x29, x30, sp, llvm#16 // 16-byte Folded Spill add x29, sp, llvm#16 str x8, sp, llvm#8. // stack canary str xzr, [sp] bl __tgt_target_data_update_nowait_mapper __tgt_target_data_update_nowait_mapper: sub sp, sp, llvm#32 stp x29, x30, sp, llvm#16 // 16-byte Folded Spill add x29, sp, llvm#16 str x8, sp, llvm#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, llvm#16 adrp x8, .L.str.8 add x8, x8, :lo12:.L.str.8 str x8, x29, llvm#32. <==. This is the insn that erases $fp ldp x29, x30, sp, llvm#16 // 16-byte Folded Reload add sp, sp, llvm#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>

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, llvm#32 stp x29, x30, sp, llvm#16 // 16-byte Folded Spill add x29, sp, llvm#16 str x8, sp, llvm#8. // stack canary str xzr, [sp] bl __tgt_target_data_update_nowait_mapper __tgt_target_data_update_nowait_mapper: sub sp, sp, llvm#32 stp x29, x30, sp, llvm#16 // 16-byte Folded Spill add x29, sp, llvm#16 str x8, sp, llvm#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, llvm#16 adrp x8, .L.str.8 add x8, x8, :lo12:.L.str.8 str x8, x29, llvm#32. <==. This is the insn that erases $fp ldp x29, x30, sp, llvm#16 // 16-byte Folded Reload add sp, sp, llvm#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>

sushgokh requested review from madhur13490, jdoerfert and jhuber6 June 22, 2024 11:50

llvmbot added the offload label Jun 22, 2024

jhuber6 reviewed Jun 22, 2024

View reviewed changes

sushgokh force-pushed the GRCO-601 branch from c96c6a1 to af4dc96 Compare June 27, 2024 07:36

llvmbot added clang Clang issues not falling into any other category clang:codegen flang:openmp clang:openmp OpenMP related changes to Clang labels Jun 27, 2024

shiltian reviewed Jun 27, 2024

View reviewed changes

clang/lib/CodeGen/CGOpenMPRuntime.cpp Outdated Show resolved Hide resolved

shiltian reviewed Jun 27, 2024

View reviewed changes

sushgokh force-pushed the GRCO-601 branch from af4dc96 to 7c6e2e6 Compare July 2, 2024 08:30

llvmbot added the llvm:transforms label Jul 2, 2024

shiltian approved these changes Jul 2, 2024

View reviewed changes

sushgokh force-pushed the GRCO-601 branch from 7c6e2e6 to b376f84 Compare July 3, 2024 10:25

sushgokh merged commit c7ee204 into llvm:main Jul 5, 2024
7 checks passed

sushgokh deleted the GRCO-601 branch July 5, 2024 05:11

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[OpenMP] Fix stack corruption due to argument mismatch #96386

[OpenMP] Fix stack corruption due to argument mismatch #96386

sushgokh commented Jun 22, 2024

llvmbot commented Jun 22, 2024 •

edited

Loading

jhuber6 Jun 22, 2024

shiltian Jun 22, 2024

sushgokh Jun 25, 2024

shiltian Jun 25, 2024

shiltian left a comment

shiltian left a comment

[OpenMP] Fix stack corruption due to argument mismatch #96386

[OpenMP] Fix stack corruption due to argument mismatch #96386

Conversation

sushgokh commented Jun 22, 2024

llvmbot commented Jun 22, 2024 • edited Loading

jhuber6 Jun 22, 2024

Choose a reason for hiding this comment

shiltian Jun 22, 2024

Choose a reason for hiding this comment

sushgokh Jun 25, 2024

Choose a reason for hiding this comment

shiltian Jun 25, 2024

Choose a reason for hiding this comment

shiltian left a comment

Choose a reason for hiding this comment

shiltian left a comment

Choose a reason for hiding this comment

llvmbot commented Jun 22, 2024 •

edited

Loading