diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -429,16 +429,16 @@
   let Category = DocCatStmt;
   let Content = [{
 If a statement is marked ``nomerge`` and contains call expressions, those call
-expressions inside the statement will not be merged during optimization. This 
+expressions inside the statement will not be merged during optimization. This
 attribute can be used to prevent the optimizer from obscuring the source
 location of certain calls. For example, it will prevent tail merging otherwise
 identical code sequences that raise an exception or terminate the program. Tail
 merging normally reduces the precision of source location information, making
 stack traces less useful for debugging. This attribute gives the user control
-over the tradeoff between code size and debug information precision. 
+over the tradeoff between code size and debug information precision.
 
-``nomerge`` attribute can also be used as function attribute to prevent all 
-calls to the specified function from merging. It has no effect on indirect 
+``nomerge`` attribute can also be used as function attribute to prevent all
+calls to the specified function from merging. It has no effect on indirect
 calls.
   }];
 }
@@ -3494,7 +3494,7 @@
 with pointers in the C family of languages. The various nullability attributes
 indicate whether a particular pointer can be null or not, which makes APIs more
 expressive and can help static analysis tools identify bugs involving null
-pointers. Clang supports several kinds of nullability attributes: the 
+pointers. Clang supports several kinds of nullability attributes: the
 ``nonnull`` and ``returns_nonnull`` attributes indicate which function or
 method parameters and result types can never be null, while nullability type
 qualifiers indicate which pointer types can be null (``_Nullable``) or cannot
@@ -3660,7 +3660,7 @@
 The ``returns_nonnull`` attribute implies that returning a null pointer is
 undefined behavior, which the optimizer may take advantage of. The ``_Nonnull``
 type qualifier indicates that a pointer cannot be null in a more general manner
-(because it is part of the type system) and does not imply undefined behavior, 
+(because it is part of the type system) and does not imply undefined behavior,
 making it more widely applicable
 }];
 }
@@ -5215,14 +5215,40 @@
 The ``callback`` attribute specifies that the annotated function may invoke the
 specified callback zero or more times. The callback, as well as the passed
 arguments, are identified by their parameter name or position (starting with
-1!) in the annotated function. The first position in the attribute identifies
-the callback callee, the following positions declare describe its arguments.
+1!) in the annotated function.
+
+The first position in the attribute identifies the callback argument encoding
+mode, ``0`` for flat mode, and ``1`` for stacked mode. The following example
+shows the difference between the two modes:
+
+  .. code-block:: c
+
+    // Flat mode
+    void flat_mode_broker(void (*cb)(int *, float *), int *arg1, float *arg2) {
+      // ...
+      cb(arg1, arg2);
+      // ...
+    }
+
+    // Stack mode
+    void stack_mode_broker(void (*cb)(int *, float *), void **args) {
+      // ...
+      cb((int *)args[0], (float *)args[1]);
+      // ...
+    }
+
+The second position in the attribute represents the callback callee, the
+following positions depend on the argument encoding mode:
+- Flat mode: describe its arguments;
+- Stack mode: describe the argument pointer. Therefore, for this mode, it can
+  only have one position.
+
 The callback callee is required to be callable with the number, and order, of
 the specified arguments. The index ``0``, or the identifier ``this``, is used to
 represent an implicit "this" pointer in class methods. If there is no implicit
 "this" pointer it shall not be referenced. The index '-1', or the name "__",
-represents an unknown callback callee argument. This can be a value which is
-not present in the declared parameter list, or one that is, but is potentially
+represents an unknown callback callee argument. This can be a value which is not
+present in the declared parameter list, or one that is, but is potentially
 inspected, captured, or modified. Parameter names and indices can be mixed in
 the callback attribute.
 
@@ -5247,14 +5273,25 @@
 
   .. code-block:: c
 
-    __attribute__((callback (start_routine, arg)))
+    __attribute__((callback (0, start_routine, arg)))
     int pthread_create(pthread_t *thread, const pthread_attr_t *attr,
                        void *(*start_routine) (void *), void *arg);
 
-    __attribute__((callback (3, 4)))
+    __attribute__((callback (0, 3, 4)))
     int pthread_create(pthread_t *thread, const pthread_attr_t *attr,
                        void *(*start_routine) (void *), void *arg);
 
+Example encodings for the callback performed by ``stack_mode_broker`` are shown
+below.
+
+  .. code-block:: c
+
+    __attribute__((callback (1, cb, arg)))
+    void stack_mode_broker(void (*cb)(int *, float *), void **args);
+
+    __attribute__((callback (1, 1, 2)))
+    void stack_mode_broker(void (*cb)(int *, float *), void **args);
+
   }];
 }
 
@@ -5478,15 +5515,15 @@
   let Content = [{
 Code can indicate CFG checks are not wanted with the ``__declspec(guard(nocf))``
 attribute. This directs the compiler to not insert any CFG checks for the entire
-function. This approach is typically used only sparingly in specific situations 
-where the programmer has manually inserted "CFG-equivalent" protection. The 
-programmer knows that they are calling through some read-only function table 
-whose address is obtained through read-only memory references and for which the 
-index is masked to the function table limit. This approach may also be applied 
-to small wrapper functions that are not inlined and that do nothing more than 
-make a call through a function pointer. Since incorrect usage of this directive 
-can compromise the security of CFG, the programmer must be very careful using 
-the directive. Typically, this usage is limited to very small functions that 
+function. This approach is typically used only sparingly in specific situations
+where the programmer has manually inserted "CFG-equivalent" protection. The
+programmer knows that they are calling through some read-only function table
+whose address is obtained through read-only memory references and for which the
+index is masked to the function table limit. This approach may also be applied
+to small wrapper functions that are not inlined and that do nothing more than
+make a call through a function pointer. Since incorrect usage of this directive
+can compromise the security of CFG, the programmer must be very careful using
+the directive. Typically, this usage is limited to very small functions that
 only call one function.
 
 `Control Flow Guard documentation <https://docs.microsoft.com/en-us/windows/win32/secbp/pe-metadata>`
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -2936,6 +2936,8 @@
   "string">;
 def err_callback_attribute_no_callee : Error<
   "'callback' attribute specifies no callback callee">;
+def err_callback_attribute_no_encoding_mode : Error<
+  "'callback' attribute specifies no parameter encoding mode">;
 def err_callback_attribute_invalid_callee : Error<
   "'callback' attribute specifies invalid callback callee">;
 def err_callback_attribute_multiple : Error<
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2121,19 +2121,23 @@
 
   if (const auto *CB = FD->getAttr<CallbackAttr>()) {
     // Annotate the callback behavior as metadata:
+    //  - The callback argument encoding mode (as argument number)
     //  - The callback callee (as argument number).
     //  - The callback payloads (as argument numbers).
     llvm::LLVMContext &Ctx = F->getContext();
     llvm::MDBuilder MDB(Ctx);
 
-    // The payload indices are all but the first one in the encoding. The first
-    // identifies the callback callee.
-    int CalleeIdx = *CB->encoding_begin();
-    ArrayRef<int> PayloadIndices(CB->encoding_begin() + 1, CB->encoding_end());
-    F->addMetadata(llvm::LLVMContext::MD_callback,
-                   *llvm::MDNode::get(Ctx, {MDB.createCallbackEncoding(
-                                               CalleeIdx, PayloadIndices,
-                                               /* VarArgsArePassed */ false)}));
+    // The payload indices are all but the first two in the encoding. The first
+    // identifies the callback argument encoding mode, and the second represents
+    // the callback callee.
+    int EncodingMode = *CB->encoding_begin();
+    int CalleeIdx = *(CB->encoding_begin() + 1);
+    ArrayRef<int> PayloadIndices(CB->encoding_begin() + 2, CB->encoding_end());
+    F->addMetadata(
+        llvm::LLVMContext::MD_callback,
+        *llvm::MDNode::get(Ctx, {MDB.createCallbackEncoding(
+                                    EncodingMode, CalleeIdx, PayloadIndices,
+                                    /* VarArgsArePassed */ false)}));
   }
 }
 
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3520,8 +3520,15 @@
 
 /// Handle __attribute__((callback(CalleeIdx, PayloadIdx0, ...))) attributes.
 static void handleCallbackAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
-  // The index that identifies the callback callee is mandatory.
+  // The parameter encoding mode is mandatory.
   if (AL.getNumArgs() == 0) {
+    S.Diag(AL.getLoc(), diag::err_callback_attribute_no_encoding_mode)
+        << AL.getRange();
+    return;
+  }
+
+  // The index that identifies the callback callee is mandatory.
+  if (AL.getNumArgs() == 1) {
     S.Diag(AL.getLoc(), diag::err_callback_attribute_no_callee)
         << AL.getRange();
     return;
@@ -3583,7 +3590,7 @@
       llvm_unreachable("Unexpected ParsedAttr argument type!");
     }
 
-    if (ArgIdx == 0 && !HasImplicitThisParam) {
+    if (ArgIdx == 0 && I != 0 && !HasImplicitThisParam) {
       S.Diag(AL.getLoc(), diag::err_callback_implicit_this_not_available)
           << (I + 1) << SR;
       return;
@@ -3597,7 +3604,9 @@
     EncodingIndices.push_back(ArgIdx);
   }
 
-  int CalleeIdx = EncodingIndices.front();
+  assert(EncodingIndices.size() >= 2);
+
+  int CalleeIdx = EncodingIndices[1];
   // Check if the callee index is proper, thus not "this" and not "unknown".
   // This means the "CalleeIdx" has to be non-negative if "HasImplicitThisParam"
   // is false and positive if "HasImplicitThisParam" is true.
@@ -3630,15 +3639,9 @@
     return;
   }
 
-  if (CalleeFnProtoType->getNumParams() > EncodingIndices.size() - 1) {
-    S.Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments)
-        << AL << (unsigned)(EncodingIndices.size() - 1);
-    return;
-  }
-
-  if (CalleeFnProtoType->getNumParams() < EncodingIndices.size() - 1) {
+  if (CalleeFnProtoType->getNumParams() != EncodingIndices.size() - 2) {
     S.Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments)
-        << AL << (unsigned)(EncodingIndices.size() - 1);
+        << AL << (unsigned)(CalleeFnProtoType->getNumParams() + 2);
     return;
   }
 
diff --git a/clang/test/CodeGen/attr-callback.c b/clang/test/CodeGen/attr-callback.c
--- a/clang/test/CodeGen/attr-callback.c
+++ b/clang/test/CodeGen/attr-callback.c
@@ -3,26 +3,26 @@
 void cb0(void);
 
 // CHECK-DAG: !callback ![[cid0:[0-9]+]] void @no_args
-__attribute__((callback(1))) void no_args(void (*callback)(void));
+__attribute__((callback(0, 1))) void no_args(void (*callback)(void));
 
 // CHECK-DAG: @args_1({{[^#]*#[0-9]+}} !callback ![[cid1:[0-9]+]]
-__attribute__((callback(1, 2, 3))) void args_1(void (*callback)(int, double), int a, double b) { no_args(cb0); }
+__attribute__((callback(0, 1, 2, 3))) void args_1(void (*callback)(int, double), int a, double b) { no_args(cb0); }
 
 // CHECK-DAG: !callback ![[cid2:[0-9]+]]  void @args_2a
-__attribute__((callback(2, 3, 3))) void args_2a(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, 2, 3, 3))) void args_2a(int a, void (*callback)(double, double), double b);
 // CHECK-DAG: !callback ![[cid2]]         void @args_2b
-__attribute__((callback(callback, b, b))) void args_2b(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, callback, b, b))) void args_2b(int a, void (*callback)(double, double), double b);
 
 // CHECK-DAG: void @args_3a({{[^#]*#[0-9]+}} !callback ![[cid3:[0-9]+]]
-__attribute__((callback(2, -1, -1))) void args_3a(int a, void (*callback)(double, double), double b) { args_2a(a, callback, b); }
+__attribute__((callback(0, 2, -1, -1))) void args_3a(int a, void (*callback)(double, double), double b) { args_2a(a, callback, b); }
 // CHECK-DAG: void @args_3b({{[^#]*#[0-9]+}} !callback ![[cid3]]
-__attribute__((callback(callback, __, __))) void args_3b(int a, void (*callback)(double, double), double b) { args_2b(a, callback, b); }
+__attribute__((callback(0, callback, __, __))) void args_3b(int a, void (*callback)(double, double), double b) { args_2b(a, callback, b); }
 
 // CHECK-DAG: ![[cid0]] = !{![[cid0b:[0-9]+]]}
-// CHECK-DAG: ![[cid0b]] = !{i64 0, i1 false}
+// CHECK-DAG: ![[cid0b]] = !{i64 0, i64 0, i1 false}
 // CHECK-DAG: ![[cid1]] = !{![[cid1b:[0-9]+]]}
-// CHECK-DAG: ![[cid1b]] = !{i64 0, i64 1, i64 2, i1 false}
+// CHECK-DAG: ![[cid1b]] = !{i64 0, i64 0, i64 1, i64 2, i1 false}
 // CHECK-DAG: ![[cid2]] = !{![[cid2b:[0-9]+]]}
-// CHECK-DAG: ![[cid2b]] = !{i64 1, i64 2, i64 2, i1 false}
+// CHECK-DAG: ![[cid2b]] = !{i64 0, i64 1, i64 2, i64 2, i1 false}
 // CHECK-DAG: ![[cid3]] = !{![[cid3b:[0-9]+]]}
-// CHECK-DAG: ![[cid3b]] = !{i64 1, i64 -1, i64 -1, i1 false}
+// CHECK-DAG: ![[cid3b]] = !{i64 0, i64 1, i64 -1, i64 -1, i1 false}
diff --git a/clang/test/CodeGen/callback_annotated.c b/clang/test/CodeGen/callback_annotated.c
--- a/clang/test/CodeGen/callback_annotated.c
+++ b/clang/test/CodeGen/callback_annotated.c
@@ -1,30 +1,30 @@
 // RUN: %clang_cc1 -triple i386-unknown-unknown -fopenmp -fno-experimental-new-pass-manager %s -emit-llvm -o - -disable-llvm-optzns | FileCheck %s --check-prefix=RUN1
 
 // RUN1-DAG: @broker0({{[^#]*#[0-9]+}} !callback ![[cid0:[0-9]+]]
-__attribute__((callback(1, 2))) void *broker0(void *(*callee)(void *), void *payload) {
+__attribute__((callback(0, 1, 2))) void *broker0(void *(*callee)(void *), void *payload) {
   return callee(payload);
 }
 
 // RUN1-DAG: @broker1({{[^#]*#[0-9]+}} !callback ![[cid1:[0-9]+]]
-__attribute__((callback(callee, payload))) void *broker1(void *payload, void *(*callee)(void *)) {
+__attribute__((callback(0, callee, payload))) void *broker1(void *payload, void *(*callee)(void *)) {
   return broker0(callee, payload);
 }
 
 void *broker2(void (*callee)(void));
 
 // RUN1-DAG: declare !callback ![[cid2:[0-9]+]] i8* @broker2
-__attribute__((callback(callee))) void *broker2(void (*callee)(void));
+__attribute__((callback(0, callee))) void *broker2(void (*callee)(void));
 
 void *broker2(void (*callee)(void));
 
 // RUN1-DAG: declare !callback ![[cid3:[0-9]+]] i8* @broker3
-__attribute__((callback(4, 1, 2, c))) void *broker3(int, int, int c, int (*callee)(int, int, int), int);
+__attribute__((callback(0, 4, 1, 2, c))) void *broker3(int, int, int c, int (*callee)(int, int, int), int);
 
 // RUN1-DAG: declare !callback ![[cid4:[0-9]+]] i8* @broker4
-__attribute__((callback(4, -1, a, __))) void *broker4(int a, int, int, int (*callee)(int, int, int), int);
+__attribute__((callback(0, 4, -1, a, __))) void *broker4(int a, int, int, int (*callee)(int, int, int), int);
 
 // RUN1-DAG: declare !callback ![[cid5:[0-9]+]] i8* @broker5
-__attribute__((callback(4, d, 5, 2))) void *broker5(int, int, int, int (*callee)(int, int, int), int d);
+__attribute__((callback(0, 4, d, 5, 2))) void *broker5(int, int, int, int (*callee)(int, int, int), int d);
 
 static void *VoidPtr2VoidPtr(void *payload) {
   return payload;
@@ -44,14 +44,14 @@
 }
 
 // RUN1-DAG: ![[cid0]] = !{![[cid0b:[0-9]+]]}
-// RUN1-DAG: ![[cid0b]] = !{i64 0, i64 1, i1 false}
+// RUN1-DAG: ![[cid0b]] = !{i64 0, i64 0, i64 1, i1 false}
 // RUN1-DAG: ![[cid1]] = !{![[cid1b:[0-9]+]]}
-// RUN1-DAG: ![[cid1b]] = !{i64 1, i64 0, i1 false}
+// RUN1-DAG: ![[cid1b]] = !{i64 0, i64 1, i64 0, i1 false}
 // RUN1-DAG: ![[cid2]] = !{![[cid2b:[0-9]+]]}
-// RUN1-DAG: ![[cid2b]] = !{i64 0, i1 false}
+// RUN1-DAG: ![[cid2b]] = !{i64 0, i64 0, i1 false}
 // RUN1-DAG: ![[cid3]] = !{![[cid3b:[0-9]+]]}
-// RUN1-DAG: ![[cid3b]] = !{i64 3, i64 0, i64 1, i64 2, i1 false}
+// RUN1-DAG: ![[cid3b]] = !{i64 0, i64 3, i64 0, i64 1, i64 2, i1 false}
 // RUN1-DAG: ![[cid4]] = !{![[cid4b:[0-9]+]]}
-// RUN1-DAG: ![[cid4b]] = !{i64 3, i64 -1, i64 0, i64 -1, i1 false}
+// RUN1-DAG: ![[cid4b]] = !{i64 0, i64 3, i64 -1, i64 0, i64 -1, i1 false}
 // RUN1-DAG: ![[cid5]] = !{![[cid5b:[0-9]+]]}
-// RUN1-DAG: ![[cid5b]] = !{i64 3, i64 4, i64 4, i64 1, i1 false}
+// RUN1-DAG: ![[cid5b]] = !{i64 0, i64 3, i64 4, i64 4, i64 1, i1 false}
diff --git a/clang/test/CodeGen/callback_openmp.c b/clang/test/CodeGen/callback_openmp.c
--- a/clang/test/CodeGen/callback_openmp.c
+++ b/clang/test/CodeGen/callback_openmp.c
@@ -3,7 +3,7 @@
 // CHECK: declare !callback ![[cid:[0-9]+]] void @__kmpc_fork_call
 // CHECK: declare !callback ![[cid]] void @__kmpc_fork_teams
 // CHECK: ![[cid]] = !{![[cidb:[0-9]+]]}
-// CHECK: ![[cidb]] = !{i64 2, i64 -1, i64 -1, i1 true}
+// CHECK: ![[cidb]] = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 
 void work1(int, int);
 void work2(int, int);
diff --git a/clang/test/CodeGenCXX/attr-callback.cpp b/clang/test/CodeGenCXX/attr-callback.cpp
--- a/clang/test/CodeGenCXX/attr-callback.cpp
+++ b/clang/test/CodeGenCXX/attr-callback.cpp
@@ -3,28 +3,28 @@
 struct Base {
 
   void no_args_1(void (*callback)(void));
-  __attribute__((callback(1))) void no_args_2(void (*callback1)(void), void (*callback2)(void));
-  __attribute__((callback(callback1))) void no_args_3(void (*callback1)(void), void (*callback2)(void));
+  __attribute__((callback(0, 1))) void no_args_2(void (*callback1)(void), void (*callback2)(void));
+  __attribute__((callback(0, callback1))) void no_args_3(void (*callback1)(void), void (*callback2)(void));
 
   // TODO: There should probably be a warning or even an error for different
   //       callbacks on the same method.
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_1(void (*callback)(void));
 
-  __attribute__((callback(callback, this, __, this))) virtual void
+  __attribute__((callback(0, callback, this, __, this))) virtual void
   this_unknown_this(void (*callback)(Base *, Base *, Base *));
 };
 
 // CHECK-DAG:      define{{.*}} void @_ZN4Base9no_args_1EPFvvE({{[^!]*!callback}} ![[cid0:[0-9]+]]
-__attribute__((callback(1))) void
+__attribute__((callback(0, 1))) void
 Base::no_args_1(void (*callback)(void)) {
 }
 
 // CHECK-DAG:      define{{.*}} void @_ZN4Base9no_args_2EPFvvES1_({{[^!]*!callback}} ![[cid1:[0-9]+]]
-__attribute__((callback(2))) void Base::no_args_2(void (*callback1)(void), void (*callback2)(void)) {
+__attribute__((callback(0, 2))) void Base::no_args_2(void (*callback1)(void), void (*callback2)(void)) {
 }
 // CHECK-DAG:      define{{.*}} void @_ZN4Base9no_args_3EPFvvES1_({{[^!]*!callback}} ![[cid1]]
-__attribute__((callback(callback2))) void Base::no_args_3(void (*callback1)(void), void (*callback2)(void)) {
+__attribute__((callback(0, callback2))) void Base::no_args_3(void (*callback1)(void), void (*callback2)(void)) {
 }
 
 // CHECK-DAG:      define{{.*}} void @_ZN4Base17this_unknown_thisEPFvPS_S0_S0_E({{[^!]*!callback}} ![[cid2:[0-9]+]]
@@ -32,7 +32,7 @@
 }
 
 struct Derived_1 : public Base {
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_1(void (*callback)(void)) override;
 };
 
@@ -48,8 +48,8 @@
 void Derived_2::virtual_1(void (*callback)(void)) {}
 
 // CHECK-DAG: ![[cid0]] = !{![[cid0b:[0-9]+]]}
-// CHECK-DAG: ![[cid0b]] = !{i64 1, i1 false}
+// CHECK-DAG: ![[cid0b]] = !{i64 0, i64 1, i1 false}
 // CHECK-DAG: ![[cid1]] = !{![[cid1b:[0-9]+]]}
-// CHECK-DAG: ![[cid1b]] = !{i64 2, i1 false}
+// CHECK-DAG: ![[cid1b]] = !{i64 0, i64 2, i1 false}
 // CHECK-DAG: ![[cid2]] = !{![[cid2b:[0-9]+]]}
-// CHECK-DAG: ![[cid2b]] = !{i64 1, i64 0, i64 -1, i64 0, i1 false}
+// CHECK-DAG: ![[cid2b]] = !{i64 0, i64 1, i64 0, i64 -1, i64 0, i1 false}
diff --git a/clang/test/OpenMP/parallel_codegen.cpp b/clang/test/OpenMP/parallel_codegen.cpp
--- a/clang/test/OpenMP/parallel_codegen.cpp
+++ b/clang/test/OpenMP/parallel_codegen.cpp
@@ -182,5 +182,5 @@
 // ALL: attributes #[[FN_ATTRS]] = {{.+}} nounwind
 // ALL-DEBUG: attributes #[[FN_ATTRS]] = {{.+}} nounwind
 // ALL: ![[cbid]] = !{![[cbidb:[0-9]+]]}
-// ALL: ![[cbidb]] = !{i64 2, i64 -1, i64 -1, i1 true}
+// ALL: ![[cbidb]] = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 #endif
diff --git a/clang/test/Sema/attr-callback-broken.c b/clang/test/Sema/attr-callback-broken.c
--- a/clang/test/Sema/attr-callback-broken.c
+++ b/clang/test/Sema/attr-callback-broken.c
@@ -1,75 +1,76 @@
 // RUN: %clang_cc1 %s -verify -fsyntax-only
 
-__attribute__((callback())) void no_callee(void (*callback)(void)); // expected-error {{'callback' attribute specifies no callback callee}}
+__attribute__((callback())) void no_encoding_mode(void (*callback)(void)); // expected-error {{'callback' attribute specifies no parameter encoding mode}}
+__attribute__((callback(0))) void no_callee(void (*callback)(void)); // expected-error {{'callback' attribute specifies no callback callee}}
 
-__attribute__((callback(1, 1))) void too_many_args_1(void (*callback)(void)) {}      // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(1, -1))) void too_many_args_2(double (*callback)(void));     // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(1, 2, 2))) void too_many_args_3(void (*callback)(int), int); // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, 1, 1))) void too_many_args_1(void (*callback)(void)) {}      // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, 1, -1))) void too_many_args_2(double (*callback)(void));     // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, 1, 2, 2))) void too_many_args_3(void (*callback)(int), int); // expected-error {{'callback' attribute requires exactly 3 arguments}}
 
-__attribute__((callback(1, 2))) void too_few_args_1(void (*callback)(int, int), int); // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(1))) void too_few_args_2(int (*callback)(int));               // expected-error {{'callback' attribute takes no arguments}}
-__attribute__((callback(1, -1))) void too_few_args_3(void (*callback)(int, int)) {}   // expected-error {{'callback' attribute takes one argument}}
+__attribute__((callback(0, 1, 2))) void too_few_args_1(void (*callback)(int, int), int); // expected-error {{'callback' attribute requires exactly 4 arguments}}
+__attribute__((callback(0, 1))) void too_few_args_2(int (*callback)(int));               // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, 1, -1))) void too_few_args_3(void (*callback)(int, int)) {}   // expected-error {{'callback' attribute requires exactly 4 arguments}}
 
-__attribute__((callback(-1))) void oob_args_1(void (*callback)(void));         // expected-error {{'callback' attribute specifies invalid callback callee}}
-__attribute__((callback(2))) void oob_args_2(int *(*callback)(void)) {}        // expected-error {{'callback' attribute parameter 1 is out of bounds}}
-__attribute__((callback(1, 3))) void oob_args_3(short (*callback)(int), int);  // expected-error {{'callback' attribute parameter 2 is out of bounds}}
-__attribute__((callback(-2, 2))) void oob_args_4(void *(*callback)(int), int); // expected-error {{'callback' attribute parameter 1 is out of bounds}}
-__attribute__((callback(1, -2))) void oob_args_5(void *(*callback)(int), int); // expected-error {{'callback' attribute parameter 2 is out of bounds}}
-__attribute__((callback(1, 2))) void oob_args_6(void *(*callback)(int), ...);  // expected-error {{'callback' attribute parameter 2 is out of bounds}}
+__attribute__((callback(0, -1))) void oob_args_1(void (*callback)(void));         // expected-error {{'callback' attribute specifies invalid callback callee}}
+__attribute__((callback(0, 2))) void oob_args_2(int *(*callback)(void)) {}        // expected-error {{'callback' attribute parameter 2 is out of bounds}}
+__attribute__((callback(0, 1, 3))) void oob_args_3(short (*callback)(int), int);  // expected-error {{'callback' attribute parameter 3 is out of bounds}}
+__attribute__((callback(0, -2, 2))) void oob_args_4(void *(*callback)(int), int); // expected-error {{'callback' attribute parameter 2 is out of bounds}}
+__attribute__((callback(0, 1, -2))) void oob_args_5(void *(*callback)(int), int); // expected-error {{'callback' attribute parameter 3 is out of bounds}}
+__attribute__((callback(0, 1, 2))) void oob_args_6(void *(*callback)(int), ...);  // expected-error {{'callback' attribute parameter 3 is out of bounds}}
 
-__attribute__((callback(1))) __attribute__((callback(1))) void multiple_cb_1(void (*callback)(void));                           // expected-error {{multiple 'callback' attributes specified}}
-__attribute__((callback(1))) __attribute__((callback(2))) void multiple_cb_2(void (*callback1)(void), void (*callback2)(void)); // expected-error {{multiple 'callback' attributes specified}}
+__attribute__((callback(0, 1))) __attribute__((callback(0, 1))) void multiple_cb_1(void (*callback)(void));                           // expected-error {{multiple 'callback' attributes specified}}
+__attribute__((callback(0, 1))) __attribute__((callback(0, 2))) void multiple_cb_2(void (*callback1)(void), void (*callback2)(void)); // expected-error {{multiple 'callback' attributes specified}}
 
 #ifdef HAS_THIS
-__attribute__((callback(0))) void oob_args_0(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
+__attribute__((callback(0, 0))) void oob_args_0(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
 #else
-__attribute__((callback(0))) void oob_args_0(void (*callback)(void));                 // expected-error {{'callback' argument at position 1 references unavailable implicit 'this'}}
-__attribute__((callback(1, 0))) void no_this_1(void *(*callback)(void *));            // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
-__attribute__((callback(1, 0))) void no_this_2(void *(*callback)(int, void *));       // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
+__attribute__((callback(0, 0))) void oob_args_0(void (*callback)(void));                 // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
+__attribute__((callback(0, 1, 0))) void no_this_1(void *(*callback)(void *));            // expected-error {{'callback' argument at position 3 references unavailable implicit 'this'}}
+__attribute__((callback(0, 1, 0))) void no_this_2(void *(*callback)(int, void *));       // expected-error {{'callback' argument at position 3 references unavailable implicit 'this'}}
 #endif
 
 // We could allow the following declarations if we at some point need to:
 
-__attribute__((callback(1, -1))) void vararg_cb_1(void (*callback)(int, ...)) {}     // expected-error {{'callback' attribute callee may not be variadic}}
-__attribute__((callback(1, 1))) void vararg_cb_2(void (*callback)(int, ...), int a); // expected-error {{'callback' attribute callee may not be variadic}}
+__attribute__((callback(0, 1, -1))) void vararg_cb_1(void (*callback)(int, ...)) {}     // expected-error {{'callback' attribute callee may not be variadic}}
+__attribute__((callback(0, 1, 1))) void vararg_cb_2(void (*callback)(int, ...), int a); // expected-error {{'callback' attribute callee may not be variadic}}
 
-__attribute__((callback(1, -1, 1, 2, 3, 4, -1))) void varargs_1(void (*callback)(int, ...), int a, float b, double c) {}               // expected-error {{'callback' attribute requires exactly 6 arguments}}
-__attribute__((callback(1, -1, 4, 2, 3, 4, -1))) void varargs_2(void (*callback)(void *, double, int, ...), int a, float b, double c); // expected-error {{'callback' attribute requires exactly 6 arguments}}
+__attribute__((callback(0, 1, -1, 1, 2, 3, 4, -1))) void varargs_1(void (*callback)(int, ...), int a, float b, double c) {}               // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, 1, -1, 4, 2, 3, 4, -1))) void varargs_2(void (*callback)(void *, double, int, ...), int a, float b, double c); // expected-error {{'callback' attribute requires exactly 5 arguments}}
 
-__attribute__((callback(1, -1, 1))) void self_arg_1(void (*callback)(int, ...)) {}          // expected-error {{'callback' attribute requires exactly 2 arguments}}
-__attribute__((callback(1, -1, 1, -1, -1, 1))) void self_arg_2(void (*callback)(int, ...)); // expected-error {{'callback' attribute requires exactly 5 arguments}}
+__attribute__((callback(0, 1, -1, 1))) void self_arg_1(void (*callback)(int, ...)) {}          // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, 1, -1, 1, -1, -1, 1))) void self_arg_2(void (*callback)(int, ...)); // expected-error {{'callback' attribute requires exactly 3 arguments}}
 
-__attribute__((callback(cb))) void unknown_name1(void (*callback)(void)) {}     // expected-error {{'callback' attribute argument 'cb' is not a known function parameter}}
-__attribute__((callback(cb, ab))) void unknown_name2(void (*cb)(int), int a) {} // expected-error {{'callback' attribute argument 'ab' is not a known function parameter}}
+__attribute__((callback(0, cb))) void unknown_name1(void (*callback)(void)) {}     // expected-error {{'callback' attribute argument 'cb' is not a known function parameter}}
+__attribute__((callback(0, cb, ab))) void unknown_name2(void (*cb)(int), int a) {} // expected-error {{'callback' attribute argument 'ab' is not a known function parameter}}
 
-__attribute__((callback(callback, 1))) void too_many_args_1b(void (*callback)(void)) {}      // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(callback, __))) void too_many_args_2b(double (*callback)(void));     // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(callback, 2, 2))) void too_many_args_3b(void (*callback)(int), int); // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, callback, 1))) void too_many_args_1b(void (*callback)(void)) {}      // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, callback, __))) void too_many_args_2b(double (*callback)(void));     // expected-error {{'callback' attribute requires exactly 2 arguments}}
+__attribute__((callback(0, callback, 2, 2))) void too_many_args_3b(void (*callback)(int), int); // expected-error {{'callback' attribute requires exactly 3 arguments}}
 
-__attribute__((callback(callback, a))) void too_few_args_1b(void (*callback)(int, int), int a); // expected-error {{'callback' attribute takes one argument}}
-__attribute__((callback(callback))) void too_few_args_2b(int (*callback)(int));                 // expected-error {{'callback' attribute takes no arguments}}
-__attribute__((callback(callback, __))) void too_few_args_3b(void (*callback)(int, int)) {}     // expected-error {{'callback' attribute takes one argument}}
+__attribute__((callback(0, callback, a))) void too_few_args_1b(void (*callback)(int, int), int a); // expected-error {{'callback' attribute requires exactly 4 arguments}}
+__attribute__((callback(0, callback))) void too_few_args_2b(int (*callback)(int));                 // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, callback, __))) void too_few_args_3b(void (*callback)(int, int)) {}     // expected-error {{'callback' attribute requires exactly 4 arguments}}
 
-__attribute__((callback(__))) void oob_args_1b(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
+__attribute__((callback(0, __))) void oob_args_1b(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
 
-__attribute__((callback(callback))) __attribute__((callback(callback))) void multiple_cb_1b(void (*callback)(void));                     // expected-error {{multiple 'callback' attributes specified}}
-__attribute__((callback(1))) __attribute__((callback(callback2))) void multiple_cb_2b(void (*callback1)(void), void (*callback2)(void)); // expected-error {{multiple 'callback' attributes specified}}
+__attribute__((callback(0, callback))) __attribute__((callback(0, callback))) void multiple_cb_1b(void (*callback)(void));                     // expected-error {{multiple 'callback' attributes specified}}
+__attribute__((callback(0, 1))) __attribute__((callback(0, callback2))) void multiple_cb_2b(void (*callback1)(void), void (*callback2)(void)); // expected-error {{multiple 'callback' attributes specified}}
 
 #ifdef HAS_THIS
-__attribute__((callback(this))) void oob_args_0b(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
+__attribute__((callback(0, this))) void oob_args_0b(void (*callback)(void)); // expected-error {{'callback' attribute specifies invalid callback callee}}
 #else
-__attribute__((callback(this))) void oob_args_0b(void (*callback)(void));           // expected-error {{'callback' argument at position 1 references unavailable implicit 'this'}}
-__attribute__((callback(1, this))) void no_this_1b(void *(*callback)(void *));      // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
-__attribute__((callback(1, this))) void no_this_2b(void *(*callback)(int, void *)); // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
+__attribute__((callback(0, this))) void oob_args_0b(void (*callback)(void));           // expected-error {{'callback' argument at position 2 references unavailable implicit 'this'}}
+__attribute__((callback(0, 1, this))) void no_this_1b(void *(*callback)(void *));      // expected-error {{'callback' argument at position 3 references unavailable implicit 'this'}}
+__attribute__((callback(0, 1, this))) void no_this_2b(void *(*callback)(int, void *)); // expected-error {{'callback' argument at position 3 references unavailable implicit 'this'}}
 #endif
 
 // We could allow the following declarations if we at some point need to:
 
-__attribute__((callback(callback, __))) void vararg_cb_1b(void (*callback)(int, ...)) {} // expected-error {{'callback' attribute callee may not be variadic}}
-__attribute__((callback(1, a))) void vararg_cb_2b(void (*callback)(int, ...), int a);    // expected-error {{'callback' attribute callee may not be variadic}}
+__attribute__((callback(0, callback, __))) void vararg_cb_1b(void (*callback)(int, ...)) {} // expected-error {{'callback' attribute callee may not be variadic}}
+__attribute__((callback(0, 1, a))) void vararg_cb_2b(void (*callback)(int, ...), int a);    // expected-error {{'callback' attribute callee may not be variadic}}
 
-__attribute__((callback(callback, __, callback, a, b, c, __))) void varargs_1b(void (*callback)(int, ...), int a, float b, double c) {} // expected-error {{'callback' attribute requires exactly 6 arguments}}
-__attribute__((callback(1, __, c, a, b, c, -1))) void varargs_2b(void (*callback)(void *, double, int, ...), int a, float b, double c); // expected-error {{'callback' attribute requires exactly 6 arguments}}
+__attribute__((callback(0, callback, __, callback, a, b, c, __))) void varargs_1b(void (*callback)(int, ...), int a, float b, double c) {} // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, 1, __, c, a, b, c, -1))) void varargs_2b(void (*callback)(void *, double, int, ...), int a, float b, double c); // expected-error {{'callback' attribute requires exactly 5 arguments}}
 
-__attribute__((callback(1, __, callback))) void self_arg_1b(void (*callback)(int, ...)) {}                        // expected-error {{'callback' attribute requires exactly 2 arguments}}
-__attribute__((callback(callback, __, callback, __, __, callback))) void self_arg_2b(void (*callback)(int, ...)); // expected-error {{'callback' attribute requires exactly 5 arguments}}
+__attribute__((callback(0, 1, __, callback))) void self_arg_1b(void (*callback)(int, ...)) {}                        // expected-error {{'callback' attribute requires exactly 3 arguments}}
+__attribute__((callback(0, callback, __, callback, __, __, callback))) void self_arg_2b(void (*callback)(int, ...)); // expected-error {{'callback' attribute requires exactly 3 arguments}}
diff --git a/clang/test/Sema/attr-callback.c b/clang/test/Sema/attr-callback.c
--- a/clang/test/Sema/attr-callback.c
+++ b/clang/test/Sema/attr-callback.c
@@ -2,13 +2,13 @@
 
 // expected-no-diagnostics
 
-__attribute__((callback(1))) void no_args(void (*callback)(void));
-__attribute__((callback(1, 2, 3))) void args_1(void (*callback)(int, double), int a, double b);
-__attribute__((callback(2, 3, 3))) void args_2(int a, void (*callback)(double, double), double b);
-__attribute__((callback(2, -1, -1))) void args_3(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, 1))) void no_args(void (*callback)(void));
+__attribute__((callback(0, 1, 2, 3))) void args_1(void (*callback)(int, double), int a, double b);
+__attribute__((callback(0, 2, 3, 3))) void args_2(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, 2, -1, -1))) void args_3(int a, void (*callback)(double, double), double b);
 
-__attribute__((callback(callback))) void no_argsb(void (*callback)(void));
-__attribute__((callback(callback, a, 3))) void args_1b(void (*callback)(int, double), int a, double b);
-__attribute__((callback(callback, b, b))) void args_2b(int a, void (*callback)(double, double), double b);
-__attribute__((callback(2, __, __))) void args_3b(int a, void (*callback)(double, double), double b);
-__attribute__((callback(callback, -1, __))) void args_3c(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, callback))) void no_argsb(void (*callback)(void));
+__attribute__((callback(0, callback, a, 3))) void args_1b(void (*callback)(int, double), int a, double b);
+__attribute__((callback(0, callback, b, b))) void args_2b(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, 2, __, __))) void args_3b(int a, void (*callback)(double, double), double b);
+__attribute__((callback(0, callback, -1, __))) void args_3c(int a, void (*callback)(double, double), double b);
diff --git a/clang/test/SemaCXX/attr-callback.cpp b/clang/test/SemaCXX/attr-callback.cpp
--- a/clang/test/SemaCXX/attr-callback.cpp
+++ b/clang/test/SemaCXX/attr-callback.cpp
@@ -9,26 +9,26 @@
 struct Base {
 
   void no_args_1(void (*callback)(void));
-  __attribute__((callback(1))) void no_args_2(void (*callback)(void));
-  __attribute__((callback(callback))) void no_args_3(void (*callback)(void)) {}
+  __attribute__((callback(0, 1))) void no_args_2(void (*callback)(void));
+  __attribute__((callback(0, callback))) void no_args_3(void (*callback)(void)) {}
 
-  __attribute__((callback(1, 0))) virtual void
+  __attribute__((callback(0, 1, 0))) virtual void
   this_tr(void (*callback)(Base *));
 
-  __attribute__((callback(1, this, __, this))) virtual void
+  __attribute__((callback(0, 1, this, __, this))) virtual void
   this_unknown_this(void (*callback)(Base *, Base *, Base *));
 
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_1(void (*callback)(void));
 
-  __attribute__((callback(callback))) virtual void
+  __attribute__((callback(0, callback))) virtual void
   virtual_2(void (*callback)(void));
 
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_3(void (*callback)(void));
 };
 
-__attribute__((callback(1))) void
+__attribute__((callback(0, 1))) void
 Base::no_args_1(void (*callback)(void)) {
 }
 
@@ -37,10 +37,10 @@
 
 struct Derived_1 : public Base {
 
-  __attribute__((callback(1, 0))) virtual void
+  __attribute__((callback(0, 1, 0))) virtual void
   this_tr(void (*callback)(Base *)) override;
 
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_1(void (*callback)(void)) override {}
 
   virtual void
@@ -49,7 +49,7 @@
 
 struct Derived_2 : public Base {
 
-  __attribute__((callback(callback))) virtual void
+  __attribute__((callback(0, callback))) virtual void
   virtual_1(void (*callback)(void)) override;
 
   virtual void
@@ -61,7 +61,7 @@
 
 void Derived_2::virtual_1(void (*callback)(void)) {}
 
-__attribute__((callback(1))) void
+__attribute__((callback(0, 1))) void
 Derived_2::virtual_2(void (*callback)(void)) {}
 
 void Derived_2::virtual_3(void (*callback)(void)) {}
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -5874,24 +5874,43 @@
 callback callee invocation to a different thread.
 
 The metadata is structured as follows: At the outer level, ``callback``
-metadata is a list of ``callback`` encodings. Each encoding starts with a
-constant ``i64`` which describes the argument position of the callback function
-in the call to the broker. The following elements, except the last, describe
-what arguments are passed to the callback function. Each element is again an
-``i64`` constant identifying the argument of the broker that is passed through,
-or ``i64 -1`` to indicate an unknown or inspected argument. The order in which
-they are listed has to be the same in which they are passed to the callback
-callee. The last element of the encoding is a boolean which specifies how
-variadic arguments of the broker are handled. If it is true, all variadic
-arguments of the broker are passed through to the callback function *after* the
-arguments encoded explicitly before.
+metadata is a list of ``callback`` encodings. Each encoding starts with two
+constants ``i64`` which describe the argument encoding mode and position of the
+callback function in the call to the broker respecitvely. Currently it supports
+two modes: flat mode and stacked mode. The following example shows the difference:
+.. code-block:: c
+
+    // Flat mode
+    void flat_mode_broker(void (*cb)(int *, float *), int *arg1, float *arg2) {
+      // ...
+      cb(arg1, arg2);
+      // ...
+    }
+
+    // Stack mode
+    void stack_mode_broker(void (*cb)(int *, float *), void **args, size_t sz) {
+      // ...
+      cb((int *)args[0], (float *)args[1]);
+      // ...
+    }
+
+The following elements, except the last, describe what arguments are passed to
+the callback function. Each element is again an ``i64`` constant identifying the
+argument of the broker that is passed through, or ``i64 -1`` to indicate an
+unknown or inspected argument. In stacked mode, it can only have two elements,
+one identifying the base pointer (the ``void **args`` shown above), and another
+identifying the size (the ``size_t sz``). The order in which they are listed has
+to be the same in which they are passed to the callback callee. The last element
+of the encoding is a boolean which specifies how variadic arguments of the broker
+are handled. If it is true, all variadic arguments of the broker are passed
+through to the callback function *after* the arguments encoded explicitly before.
 
 In the code below, the ``pthread_create`` function is marked as a broker
 through the ``!callback !1`` metadata. In the example, there is only one
 callback encoding, namely ``!2``, associated with the broker. This encoding
-identifies the callback function as the second argument of the broker (``i64
-2``) and the sole argument of the callback function as the third one of the
-broker function (``i64 3``).
+identifies the callback argument encoding mode as flat mode (``i64 0``), and
+callback function as the second argument of the broker (``i64 2``) and the sole
+argument of the callback function as the third one of the broker function (``i64 3``).
 
 .. FIXME why does the llvm-sphinx-docs builder give a highlighting
    error if the below is set to highlight as 'llvm', despite that we
@@ -5902,14 +5921,14 @@
     declare !callback !1 dso_local i32 @pthread_create(i64*, %union.pthread_attr_t*, i8* (i8*)*, i8*)
 
     ...
-    !2 = !{i64 2, i64 3, i1 false}
+    !2 = !{i64 0, i64 2, i64 3, i1 false}
     !1 = !{!2}
 
-Another example is shown below. The callback callee is the second argument of
-the ``__kmpc_fork_call`` function (``i64 2``). The callee is given two unknown
-values (each identified by a ``i64 -1``) and afterwards all
-variadic arguments that are passed to the ``__kmpc_fork_call`` call (due to the
-final ``i1 true``).
+Another example is shown below. The argument encoding is in flat mode (``i64 0``).
+The callback callee is the second argument of the ``__kmpc_fork_call``
+function (``i64 2``). The callee is given two unknown values (each identified by
+a ``i64 -1``) and afterwards all variadic arguments that are passed to the
+``__kmpc_fork_call`` call (due to the final ``i1 true``).
 
 .. FIXME why does the llvm-sphinx-docs builder give a highlighting
    error if the below is set to highlight as 'llvm', despite that we
@@ -5920,7 +5939,20 @@
     declare !callback !0 dso_local void @__kmpc_fork_call(%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...)
 
     ...
-    !1 = !{i64 2, i64 -1, i64 -1, i1 true}
+    !1 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
+    !0 = !{!1}
+
+One more example is shown below. The argument encoding is in stacked mode
+(``i64 1``), The callback callee is the second argument of the ``broker``
+function (``i64 1``). Since it is in stacked mode, the next two elements identify
+the arguments pointer(``i64 2``).
+
+.. code-block:: text
+
+    declare !callback !0 dso_local void @broker(void *, void (i32*, i32*)*, void **)
+
+    ...
+    !1 = !{i64 0, i64 1, i64 2, i1 false}
     !0 = !{!1}
 
 
@@ -11812,7 +11844,7 @@
 
       declare token
         @llvm.experimental.gc.statepoint(i64 <id>, i32 <num patch bytes>,
-                       func_type <target>, 
+                       func_type <target>,
                        i64 <#call args>, i64 <flags>,
                        ... (call parameters),
                        i64 0, i64 0)
@@ -11922,7 +11954,7 @@
 
 The first and only argument is the ``gc.statepoint`` which starts
 the safepoint sequence of which this ``gc.result`` is a part.
-Despite the typing of this as a generic token, *only* the value defined 
+Despite the typing of this as a generic token, *only* the value defined
 by a ``gc.statepoint`` is legal here.
 
 Semantics:
@@ -11946,8 +11978,8 @@
 ::
 
       declare <pointer type>
-        @llvm.experimental.gc.relocate(token %statepoint_token, 
-                                       i32 %base_offset, 
+        @llvm.experimental.gc.relocate(token %statepoint_token,
+                                       i32 %base_offset,
                                        i32 %pointer_offset)
 
 Overview:
@@ -11961,7 +11993,7 @@
 
 The first argument is the ``gc.statepoint`` which starts the
 safepoint sequence of which this ``gc.relocation`` is a part.
-Despite the typing of this as a generic token, *only* the value defined 
+Despite the typing of this as a generic token, *only* the value defined
 by a ``gc.statepoint`` is legal here.
 
 The second and third arguments are both indices into operands of the
diff --git a/llvm/include/llvm/IR/AbstractCallSite.h b/llvm/include/llvm/IR/AbstractCallSite.h
--- a/llvm/include/llvm/IR/AbstractCallSite.h
+++ b/llvm/include/llvm/IR/AbstractCallSite.h
@@ -52,6 +52,7 @@
 
   /// The encoding of a callback with regards to the underlying instruction.
   struct CallbackInfo {
+    enum ParameterEncodingModeTy : int { FLAT = 0, STACK = 1 };
 
     /// For direct/indirect calls the parameter encoding is empty. If it is not,
     /// the abstract call site represents a callback. In that case, the first
@@ -65,7 +66,6 @@
     /// unknown values that are passed to the callee.
     using ParameterEncodingTy = SmallVector<int, 0>;
     ParameterEncodingTy ParameterEncoding;
-
   };
 
 private:
@@ -147,15 +147,15 @@
       if (CE->hasOneUse() && CE->isCast())
         U = &*CE->use_begin();
 
-    return (int)CB->getArgOperandNo(U) == CI.ParameterEncoding[0];
+    return (int)CB->getArgOperandNo(U) == CI.ParameterEncoding[1];
   }
 
   /// Return the number of parameters of the callee.
   unsigned getNumArgOperands() const {
     if (isDirectCall())
       return CB->getNumArgOperands();
-    // Subtract 1 for the callee encoding.
-    return CI.ParameterEncoding.size() - 1;
+    // Subtract 2 for the argument encoding mode and callee encoding.
+    return CI.ParameterEncoding.size() - 2;
   }
 
   /// Return the operand index of the underlying instruction associated with @p
@@ -169,8 +169,8 @@
   int getCallArgOperandNo(unsigned ArgNo) const {
     if (isDirectCall())
       return ArgNo;
-    // Add 1 for the callee encoding.
-    return CI.ParameterEncoding[ArgNo + 1];
+    // Add 2 for the argument encoding mode and callee encoding.
+    return CI.ParameterEncoding[ArgNo + 2];
   }
 
   /// Return the operand of the underlying instruction associated with @p Arg.
@@ -180,21 +180,61 @@
 
   /// Return the operand of the underlying instruction associated with the
   /// function parameter number @p ArgNo or nullptr if there is none.
-  Value *getCallArgOperand(unsigned ArgNo) const {
-    if (isDirectCall())
-      return CB->getArgOperand(ArgNo);
-    // Add 1 for the callee encoding.
-    return CI.ParameterEncoding[ArgNo + 1] >= 0
-               ? CB->getArgOperand(CI.ParameterEncoding[ArgNo + 1])
-               : nullptr;
-  }
+  /// For direct call, it directly returns the corresponding parameter.
+  /// For callback call, there are two cases based on how arguments are encoded:
+  /// 1. Flat mode: return the argument based on the parameter encoding.
+  /// 2. Stack mode: return the corresponding pointer stored in the base pointer
+  /// (as shown below).
+  ///
+  /// For example, we have the following OpenMP code:
+  ///
+  /// define dso_local void @callback(i64, i64, i64) {
+  ///   ret void
+  /// }
+  /// define dso_local void @func() {
+  ///   %a = alloca i32, align 4
+  ///   %b = alloca float, align 4
+  ///   %a.casted = alloca i64, align 8
+  ///   %b.casted = alloca i64, align 8
+  ///   %.ptrs = alloca [3 x i8*], align 8
+  ///   %0 = load i32, i32* %a, align 4
+  ///   %conv = bitcast i64* %a.casted to i32*
+  ///   store i32 %0, i32* %conv, align 4
+  ///   %1 = load i64, i64* %a.casted, align 8
+  ///   %2 = load float, float* %b, align 4
+  ///   %conv1 = bitcast i64* %b.casted to float*
+  ///   store float %2, float* %conv1, align 4
+  ///   %3 = load i64, i64* %b.casted, align 8
+  ///   %6 = getelementptr inbounds [3 x i8*], [3 x i8*]* %.ptrs, i32 0, i32 0
+  ///   %7 = bitcast i8** %6 to i64*
+  ///   store i64 %1, i64* %7, align 8
+  ///   %8 = getelementptr inbounds [3 x i8*], [3 x i8*]* %.ptrs, i32 0, i32 1
+  ///   %9 = bitcast i8** %8 to i64*
+  ///   store i64 %3, i64* %9, align 8
+  ///   %10 = getelementptr inbounds [3 x i8*], [3 x i8*]* %.ptrs, i32 0, i32 0
+  ///   ^
+  ///   ; CUDA doesn't have this instruction. Instead, it directly uses %.ptrs
+  ///   ; below.
+  ///   call i32 @__tgt_target(..., i8* @callback, i32 3, i8** null, \
+  ///                          i8** %10, ...)
+  ///   ret void
+  /// }
+  /// declare !callback !1 i32 @__tgt_target(i64, i8*, i32, i8**, i8**, i64*,
+  /// i64*) !1 = !{!2} !2 = !{i64 1, i64 1, i64 4, i1 false}
+  ///
+  /// When \p getCallArgOperand is called, it gets the argument pointer (%12)
+  /// and go through all its uses to find the getelementptr instruction with the
+  /// right index (equal to ArgNo). Next, find the bitcast and store instruction
+  /// through the use chain, and finally get the right argument pointer (if
+  /// ArgNo is 1, then %3 will be returned).
+  Value *getCallArgOperand(unsigned ArgNo) const;
 
   /// Return the operand index of the underlying instruction associated with the
   /// callee of this ACS. Only valid for callback calls!
   int getCallArgOperandNoForCallee() const {
     assert(isCallbackCall());
-    assert(CI.ParameterEncoding.size() && CI.ParameterEncoding[0] >= 0);
-    return CI.ParameterEncoding[0];
+    assert(CI.ParameterEncoding.size() >= 2 && CI.ParameterEncoding[1] >= 0);
+    return CI.ParameterEncoding[1];
   }
 
   /// Return the use of the callee value in the underlying instruction. Only
diff --git a/llvm/include/llvm/IR/MDBuilder.h b/llvm/include/llvm/IR/MDBuilder.h
--- a/llvm/include/llvm/IR/MDBuilder.h
+++ b/llvm/include/llvm/IR/MDBuilder.h
@@ -102,7 +102,8 @@
   //===------------------------------------------------------------------===//
 
   /// Return metadata describing a callback (see llvm::AbstractCallSite).
-  MDNode *createCallbackEncoding(unsigned CalleeArgNo, ArrayRef<int> Arguments,
+  MDNode *createCallbackEncoding(unsigned EncodeMode, unsigned CalleeArgNo,
+                                 ArrayRef<int> Arguments,
                                  bool VarArgsArePassed);
 
   /// Merge the new callback encoding \p NewCB into \p ExistingCallbacks.
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -94,10 +94,10 @@
         //  - The first two arguments of the callback callee are unknown (-1).
         //  - All variadic arguments to the runtime function are passed to the
         //    callback callee.
-        Fn->addMetadata(
-            LLVMContext::MD_callback,
-            *MDNode::get(Ctx, {MDB.createCallbackEncoding(
-                                  2, {-1, -1}, /* VarArgsArePassed */ true)}));
+        Fn->addMetadata(LLVMContext::MD_callback,
+                        *MDNode::get(Ctx, {MDB.createCallbackEncoding(
+                                              0, 2, {-1, -1},
+                                              /* VarArgsArePassed */ true)}));
       }
     }
 
@@ -597,7 +597,7 @@
       F->addMetadata(
           llvm::LLVMContext::MD_callback,
           *llvm::MDNode::get(
-              Ctx, {MDB.createCallbackEncoding(2, {-1, -1},
+              Ctx, {MDB.createCallbackEncoding(0, 2, {-1, -1},
                                                /* VarArgsArePassed */ true)}));
     }
   }
diff --git a/llvm/lib/IR/AbstractCallSite.cpp b/llvm/lib/IR/AbstractCallSite.cpp
--- a/llvm/lib/IR/AbstractCallSite.cpp
+++ b/llvm/lib/IR/AbstractCallSite.cpp
@@ -16,6 +16,7 @@
 
 #include "llvm/IR/AbstractCallSite.h"
 #include "llvm/ADT/Statistic.h"
+#include "llvm/IR/Instructions.h"
 #include "llvm/Support/Debug.h"
 
 using namespace llvm;
@@ -44,7 +45,7 @@
 
   for (const MDOperand &Op : CallbackMD->operands()) {
     MDNode *OpMD = cast<MDNode>(Op.get());
-    auto *CBCalleeIdxAsCM = cast<ConstantAsMetadata>(OpMD->getOperand(0));
+    auto *CBCalleeIdxAsCM = cast<ConstantAsMetadata>(OpMD->getOperand(1));
     uint64_t CBCalleeIdx =
         cast<ConstantInt>(CBCalleeIdxAsCM->getValue())->getZExtValue();
     if (CBCalleeIdx < CB.arg_size())
@@ -102,7 +103,7 @@
   MDNode *CallbackEncMD = nullptr;
   for (const MDOperand &Op : CallbackMD->operands()) {
     MDNode *OpMD = cast<MDNode>(Op.get());
-    auto *CBCalleeIdxAsCM = cast<ConstantAsMetadata>(OpMD->getOperand(0));
+    auto *CBCalleeIdxAsCM = cast<ConstantAsMetadata>(OpMD->getOperand(1));
     uint64_t CBCalleeIdx =
         cast<ConstantInt>(CBCalleeIdxAsCM->getValue())->getZExtValue();
     if (CBCalleeIdx != UseIdx)
@@ -119,10 +120,12 @@
 
   NumCallbackCallSites++;
 
-  assert(CallbackEncMD->getNumOperands() >= 2 && "Incomplete !callback metadata");
+  assert(CallbackEncMD->getNumOperands() >= 3 &&
+         "Incomplete !callback metadata");
 
   unsigned NumCallOperands = CB->getNumArgOperands();
-  // Skip the var-arg flag at the end when reading the metadata.
+  // Skip the var-arg flag at the end and mode at the beginning when reading the
+  // metadata.
   for (unsigned u = 0, e = CallbackEncMD->getNumOperands() - 1; u < e; u++) {
     Metadata *OpAsM = CallbackEncMD->getOperand(u).get();
     auto *OpAsCM = cast<ConstantAsMetadata>(OpAsM);
@@ -152,3 +155,56 @@
   for (unsigned u = Callee->arg_size(); u < NumCallOperands; u++)
     CI.ParameterEncoding.push_back(u);
 }
+
+Value *AbstractCallSite::getCallArgOperand(unsigned ArgNo) const {
+  if (isDirectCall())
+    return CB->getArgOperand(ArgNo);
+
+  switch (CI.ParameterEncoding[0]) {
+  default:
+    llvm_unreachable("Unknown encoding mode");
+  case CallbackInfo::ParameterEncodingModeTy::FLAT:
+    // Add 2 for the argument encoding mode and callee encoding.
+    return CI.ParameterEncoding[ArgNo + 2] >= 0
+               ? CB->getArgOperand(CI.ParameterEncoding[ArgNo + 2])
+               : nullptr;
+  case CallbackInfo::ParameterEncodingModeTy::STACK:
+    Value *ArgsOp = CB->getArgOperand(CI.ParameterEncoding[2]);
+    // For OpenMP, it has one more instruction to get the address of the first
+    // element of Args. CUDA doesn't do that.
+    if (ArgsOp->hasOneUse())
+      if (GetElementPtrInst *GEPInst = dyn_cast<GetElementPtrInst>(ArgsOp))
+        ArgsOp = GEPInst->getOperand(0);
+
+    // Find the GetElementPtrInst with the right index
+    Value *GEP = nullptr;
+    for (Use &U : ArgsOp->uses()) {
+      if (GetElementPtrInst *GEPInst =
+              dyn_cast<GetElementPtrInst>(U.getUser())) {
+        Value *Idx = GEPInst->getOperand(GEPInst->getNumOperands() - 1);
+        ConstantInt *CI = dyn_cast<ConstantInt>(Idx);
+        if (!CI)
+          continue;
+
+        if (CI->getZExtValue() == ArgNo) {
+          GEP = U.getUser();
+          break;
+        }
+      }
+    }
+
+    if (!GEP || !GEP->hasOneUse())
+      return nullptr;
+
+    if (BitCastInst *CE = dyn_cast<BitCastInst>(GEP->use_begin()->getUser()))
+      if (CE->hasOneUse())
+        GEP = CE;
+
+    // The only use of the bitcast instruction is to store the argument to Args.
+    StoreInst *SI = dyn_cast<StoreInst>(GEP->use_begin()->getUser());
+    if (!SI)
+      return nullptr;
+
+    return SI->getOperand(0);
+  }
+}
diff --git a/llvm/lib/IR/MDBuilder.cpp b/llvm/lib/IR/MDBuilder.cpp
--- a/llvm/lib/IR/MDBuilder.cpp
+++ b/llvm/lib/IR/MDBuilder.cpp
@@ -104,12 +104,14 @@
   return MDNode::get(Context, Ops);
 }
 
-MDNode *MDBuilder::createCallbackEncoding(unsigned CalleeArgNo,
+MDNode *MDBuilder::createCallbackEncoding(unsigned EncodeMode,
+                                          unsigned CalleeArgNo,
                                           ArrayRef<int> Arguments,
                                           bool VarArgArePassed) {
   SmallVector<Metadata *, 4> Ops;
 
   Type *Int64 = Type::getInt64Ty(Context);
+  Ops.push_back(createConstant(ConstantInt::get(Int64, EncodeMode)));
   Ops.push_back(createConstant(ConstantInt::get(Int64, CalleeArgNo)));
 
   for (int ArgNo : Arguments)
diff --git a/llvm/test/Analysis/CallGraph/callback-calls.ll b/llvm/test/Analysis/CallGraph/callback-calls.ll
--- a/llvm/test/Analysis/CallGraph/callback-calls.ll
+++ b/llvm/test/Analysis/CallGraph/callback-calls.ll
@@ -17,4 +17,4 @@
 declare !callback !0 void @broker(void (i32*)*, i32*)
 
 !0 = !{!1}
-!1 = !{i64 0, i64 1, i1 false}
+!1 = !{i64 0, i64 0, i64 1, i1 false}
diff --git a/llvm/test/Analysis/CallGraph/ignore-callback-uses.ll b/llvm/test/Analysis/CallGraph/ignore-callback-uses.ll
--- a/llvm/test/Analysis/CallGraph/ignore-callback-uses.ll
+++ b/llvm/test/Analysis/CallGraph/ignore-callback-uses.ll
@@ -48,4 +48,4 @@
 declare !callback !2 void @__kmpc_fork_call(%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) #2
 
 !2 = !{!3}
-!3 = !{i64 2, i64 -1, i64 -1, i1 true}
+!3 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
diff --git a/llvm/test/Transforms/Attributor/IPConstantProp/multiple_callbacks.ll b/llvm/test/Transforms/Attributor/IPConstantProp/multiple_callbacks.ll
--- a/llvm/test/Transforms/Attributor/IPConstantProp/multiple_callbacks.ll
+++ b/llvm/test/Transforms/Attributor/IPConstantProp/multiple_callbacks.ll
@@ -146,9 +146,9 @@
 
 declare !callback !3 void @broker(i32 (i32)*, i32 (i32)*, i32 (i32)*, i32, i32)
 
-!0 = !{i64 0, i64 3, i1 false}
-!1 = !{i64 1, i64 4, i1 false}
-!2 = !{i64 2, i64 3, i1 false}
+!0 = !{i64 0, i64 0, i64 3, i1 false}
+!1 = !{i64 0, i64 1, i64 4, i1 false}
+!2 = !{i64 0, i64 2, i64 3, i1 false}
 !3 = !{!0, !2, !1}
 ;.
 ; IS__TUNIT____: attributes #[[ATTR0]] = { nofree nosync nounwind readnone willreturn }
@@ -156,7 +156,7 @@
 ; IS__CGSCC____: attributes #[[ATTR0]] = { nofree norecurse nosync nounwind readnone willreturn }
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1, !2, !3}
-; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 3, i1 false}
-; CHECK: [[META2:![0-9]+]] = !{i64 2, i64 3, i1 false}
-; CHECK: [[META3:![0-9]+]] = !{i64 1, i64 4, i1 false}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 0, i64 3, i1 false}
+; CHECK: [[META2:![0-9]+]] = !{i64 0, i64 2, i64 3, i1 false}
+; CHECK: [[META3:![0-9]+]] = !{i64 0, i64 1, i64 4, i1 false}
 ;.
diff --git a/llvm/test/Transforms/Attributor/IPConstantProp/openmp_parallel_for.ll b/llvm/test/Transforms/Attributor/IPConstantProp/openmp_parallel_for.ll
--- a/llvm/test/Transforms/Attributor/IPConstantProp/openmp_parallel_for.ll
+++ b/llvm/test/Transforms/Attributor/IPConstantProp/openmp_parallel_for.ll
@@ -290,9 +290,9 @@
 
 declare !callback !0 dso_local void @__kmpc_fork_call(%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...)
 
-!1 = !{i64 2, i64 -1, i64 -1, i1 true}
+!1 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 !0 = !{!1}
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1}
-; CHECK: [[META1:![0-9]+]] = !{i64 2, i64 -1, i64 -1, i1 true}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 ;.
diff --git a/llvm/test/Transforms/Attributor/IPConstantProp/pthreads.ll b/llvm/test/Transforms/Attributor/IPConstantProp/pthreads.ll
--- a/llvm/test/Transforms/Attributor/IPConstantProp/pthreads.ll
+++ b/llvm/test/Transforms/Attributor/IPConstantProp/pthreads.ll
@@ -137,7 +137,7 @@
   ret i8* %arg
 }
 
-!1 = !{i64 2, i64 3, i1 false}
+!1 = !{i64 0, i64 2, i64 3, i1 false}
 !0 = !{!1}
 ;.
 ; IS__TUNIT____: attributes #[[ATTR0]] = { nofree nosync nounwind readnone willreturn }
@@ -145,5 +145,5 @@
 ; IS__CGSCC____: attributes #[[ATTR0]] = { nofree norecurse nosync nounwind readnone willreturn }
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1}
-; CHECK: [[META1:![0-9]+]] = !{i64 2, i64 3, i1 false}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 2, i64 3, i1 false}
 ;.
diff --git a/llvm/test/Transforms/Attributor/IPConstantProp/thread_local_acs.ll b/llvm/test/Transforms/Attributor/IPConstantProp/thread_local_acs.ll
--- a/llvm/test/Transforms/Attributor/IPConstantProp/thread_local_acs.ll
+++ b/llvm/test/Transforms/Attributor/IPConstantProp/thread_local_acs.ll
@@ -71,7 +71,7 @@
 
 declare !callback !0 dso_local void @broker(i32*, i32 (i32*, i32*)*, i32*)
 
-!1 = !{i64 1, i64 0, i64 2, i1 false}
+!1 = !{i64 0, i64 1, i64 0, i64 2, i1 false}
 !0 = !{!1}
 ;.
 ; IS__TUNIT____: attributes #[[ATTR0]] = { nofree nosync nounwind readonly willreturn }
@@ -79,5 +79,5 @@
 ; IS__CGSCC____: attributes #[[ATTR0]] = { nofree norecurse nosync nounwind readonly willreturn }
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1}
-; CHECK: [[META1:![0-9]+]] = !{i64 1, i64 0, i64 2, i1 false}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 1, i64 0, i64 2, i1 false}
 ;.
diff --git a/llvm/test/Transforms/Attributor/callbacks.ll b/llvm/test/Transforms/Attributor/callbacks.ll
--- a/llvm/test/Transforms/Attributor/callbacks.ll
+++ b/llvm/test/Transforms/Attributor/callbacks.ll
@@ -417,10 +417,10 @@
 declare !callback !0 void @t3_callback_broker(i32* nocapture , i32* nocapture , void (i32*, i32*, ...)* nocapture, ...)
 
 !0 = !{!1}
-!1 = !{i64 2, i64 -1, i64 -1, i1 true}
+!1 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 ;.
 ; CHECK: attributes #[[ATTR0:[0-9]+]] = { nosync }
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1}
-; CHECK: [[META1:![0-9]+]] = !{i64 2, i64 -1, i64 -1, i1 true}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 ;.
diff --git a/llvm/test/Transforms/Attributor/noundef.ll b/llvm/test/Transforms/Attributor/noundef.ll
--- a/llvm/test/Transforms/Attributor/noundef.ll
+++ b/llvm/test/Transforms/Attributor/noundef.ll
@@ -122,9 +122,9 @@
 }
 
 declare !callback !0 void @callback_broker(void (i8*)*, i8*)
-!1 = !{i64 0, i64 1, i1 false}
+!1 = !{i64 0, i64 0, i64 1, i1 false}
 !0 = !{!1}
 ;.
 ; CHECK: [[META0:![0-9]+]] = !{!1}
-; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 1, i1 false}
+; CHECK: [[META1:![0-9]+]] = !{i64 0, i64 0, i64 1, i1 false}
 ;.
diff --git a/llvm/test/Transforms/OpenMP/parallel_deletion.ll b/llvm/test/Transforms/OpenMP/parallel_deletion.ll
--- a/llvm/test/Transforms/OpenMP/parallel_deletion.ll
+++ b/llvm/test/Transforms/OpenMP/parallel_deletion.ll
@@ -442,7 +442,7 @@
 !0 = !{i32 1, !"wchar_size", i32 4}
 !1 = !{!"clang"}
 !2 = !{!3}
-!3 = !{i64 2, i64 -1, i64 -1, i1 true}
+!3 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 !4 = !{!5, !5, i64 0}
 !5 = !{!"int", !6, i64 0}
 !6 = !{!"omnipotent char", !7, i64 0}
diff --git a/llvm/test/Transforms/OpenMP/parallel_deletion_cg_update.ll b/llvm/test/Transforms/OpenMP/parallel_deletion_cg_update.ll
--- a/llvm/test/Transforms/OpenMP/parallel_deletion_cg_update.ll
+++ b/llvm/test/Transforms/OpenMP/parallel_deletion_cg_update.ll
@@ -9,7 +9,7 @@
 ; CHECK:   CS<None> calls function 'd'
 ;
 ; CHECK: Call graph node for function: '.omp_outlined..0'<<{{.*}}>>  #uses=1
-; 
+;
 ; CHECK: Call graph node for function: '.omp_outlined..1'<<{{.*}}>>  #uses=3
 ; CHECK:   CS<{{.*}}> calls function 'd'
 ;
@@ -89,4 +89,4 @@
 !0 = !{i32 1, !"wchar_size", i32 4}
 !1 = !{!"clang version 11.0.0"}
 !2 = !{!3}
-!3 = !{i64 2, i64 -1, i64 -1, i1 true}
+!3 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
diff --git a/llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll b/llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll
--- a/llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll
+++ b/llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll
@@ -96,7 +96,7 @@
 !21 = !DILocation(line: 14, column: 1, scope: !15)
 !22 = !DILocation(line: 16, column: 1, scope: !15)
 !23 = !{!24}
-!24 = !{i64 2, i64 -1, i64 -1, i1 true}
+!24 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 !25 = distinct !DISubprogram(name: ".omp_outlined.", scope: !1, file: !1, line: 9, type: !26, scopeLine: 9, flags: DIFlagPrototyped | DIFlagAllCallsDescribed, spFlags: DISPFlagLocalToUnit | DISPFlagDefinition | DISPFlagOptimized, unit: !0, retainedNodes: !33)
 !26 = !DISubroutineType(types: !27)
 !27 = !{null, !28, !28}
diff --git a/llvm/test/Transforms/OpenMP/parallel_region_merging.ll b/llvm/test/Transforms/OpenMP/parallel_region_merging.ll
--- a/llvm/test/Transforms/OpenMP/parallel_region_merging.ll
+++ b/llvm/test/Transforms/OpenMP/parallel_region_merging.ll
@@ -786,7 +786,7 @@
 
 !0 = !{i32 1, !"wchar_size", i32 4}
 !1 = !{!2}
-!2 = !{i64 2, i64 -1, i64 -1, i1 true}
+!2 = !{i64 0, i64 2, i64 -1, i64 -1, i1 true}
 ; CHECK-LABEL: define {{[^@]+}}@merge
 ; CHECK-SAME: (i32 [[A:%.*]]) local_unnamed_addr {
 ; CHECK-NEXT:  entry:
diff --git a/llvm/unittests/IR/AbstractCallSiteTest.cpp b/llvm/unittests/IR/AbstractCallSiteTest.cpp
--- a/llvm/unittests/IR/AbstractCallSiteTest.cpp
+++ b/llvm/unittests/IR/AbstractCallSiteTest.cpp
@@ -23,7 +23,7 @@
   return Mod;
 }
 
-TEST(AbstractCallSite, CallbackCall) {
+TEST(AbstractCallSite, FlatCallbackCall) {
   LLVMContext C;
 
   const char *IR =
@@ -36,7 +36,7 @@
       "}\n"
       "declare !callback !0 void @broker(i32, void (i8*, ...)*, ...)\n"
       "!0 = !{!1}\n"
-      "!1 = !{i64 1, i64 -1, i1 true}";
+      "!1 = !{i64 0, i64 1, i64 -1, i1 true}";
 
   std::unique_ptr<Module> M = parseIR(C, IR);
   ASSERT_TRUE(M);
@@ -53,3 +53,105 @@
   EXPECT_TRUE(ACS.isCallee(CallbackUse));
   EXPECT_EQ(ACS.getCalledFunction(), Callback);
 }
+
+TEST(AbstractCallSite, StackedCallbackCallOpenMP) {
+  LLVMContext C;
+
+  const char *IR =
+    "define dso_local void @callback(i64, i64) {\n"
+    "  ret void\n"
+    "}\n"
+    "define dso_local void @broker() {\n"
+    "entry:\n"
+    "  %a = alloca i32, align 4\n"
+    "  %b = alloca float, align 4\n"
+    "  %c = alloca double, align 8\n"
+    "  %a.casted = alloca i64, align 8\n"
+    "  %b.casted = alloca i64, align 8\n"
+    "  %c.casted = alloca i64, align 8\n"
+    "  %args = alloca [3 x i8*], align 8\n"
+    "  %0 = load i32, i32* %a, align 4\n"
+    "  %conv = bitcast i64* %a.casted to i32*\n"
+    "  store i32 %0, i32* %conv, align 4\n"
+    "  %1 = load i64, i64* %a.casted, align 8\n"
+    "  %2 = load float, float* %b, align 4\n"
+    "  %conv1 = bitcast i64* %b.casted to float*\n"
+    "  store float %2, float* %conv1, align 4\n"
+    "  %b.addr = load i64, i64* %b.casted, align 8\n"
+    "  %3 = getelementptr inbounds [3 x i8*], [3 x i8*]* %args, i32 0, i32 0\n"
+    "  %4 = bitcast i8** %3 to i64*\n"
+    "  store i64 %1, i64* %4, align 8\n"
+    "  %5 = getelementptr inbounds [3 x i8*], [3 x i8*]* %args, i32 0, i32 1\n"
+    "  %6 = bitcast i8** %5 to i64*\n"
+    "  store i64 %b.addr, i64* %6, align 8\n"
+    "  %7 = getelementptr inbounds [3 x i8*], [3 x i8*]* %args, i32 0, i32 0\n"
+    "  call i32 @__tgt_target(i64 -1, i8* bitcast (void (i64, i64)* @callback to i8*), i32 2, i8** null, i8** %7, i64* null, i64* null)\n"
+    "  ret void\n"
+    "}\n"
+    "declare !callback !1 i32 @__tgt_target(i64, i8*, i32, i8**, i8**, i64*, i64*)\n"
+    "!1 = !{!2}\n"
+    "!2 = !{i64 1, i64 1, i64 4, i1 false}\n";
+
+  std::unique_ptr<Module> M = parseIR(C, IR);
+  ASSERT_TRUE(M);
+
+  Function *Callback = M->getFunction("callback");
+  ASSERT_NE(Callback, nullptr);
+
+  const Use *CallbackUse = Callback->getSingleUndroppableUse();
+  ASSERT_NE(CallbackUse, nullptr);
+
+  AbstractCallSite ACS(CallbackUse);
+  EXPECT_TRUE(ACS);
+  EXPECT_TRUE(ACS.isCallbackCall());
+  EXPECT_TRUE(ACS.isCallee(CallbackUse));
+  EXPECT_EQ(ACS.getCalledFunction(), Callback);
+  EXPECT_TRUE(ACS.getCallArgOperand(1));
+  EXPECT_EQ(ACS.getCallArgOperand(1)->getName(), "b.addr");
+}
+
+TEST(AbstractCallSite, StackedCallbackCallCUDA) {
+  LLVMContext C;
+
+  const char *IR =
+    "%struct.CUstream_st = type opaque\n"
+    "define dso_local void @callback(i32 %n, float %a) {\n"
+    "  ret void\n"
+    "}\n"
+    "define dso_local void @broker(i32 %n, float %a) {\n"
+    "entry:\n"
+    "  %n.addr = alloca i32, align 4\n"
+    "  %a.addr = alloca float, align 4\n"
+    "  store i32 %n, i32* %n.addr, align 4\n"
+    "  store float %a, float* %a.addr, align 4\n"
+    "  %args = alloca i8*, i64 4, align 16\n"
+    "  %0 = bitcast i32* %n.addr to i8*\n"
+    "  %1 = getelementptr i8*, i8** %args, i32 0\n"
+    "  store i8* %0, i8** %1, align 8\n"
+    "  %a.addr.cast = bitcast float* %a.addr to i8*\n"
+    "  %2 = getelementptr i8*, i8** %args, i32 1\n"
+    "  store i8* %a.addr.cast, i8** %2, align 8\n"
+    "  call i32 @cudaLaunchKernel(i8* bitcast (void (i32, float)* @callback to i8*), i64 0, i32 0, i64 0, i32 0, i8** %args, i64 0, %struct.CUstream_st* null)\n"
+    "  ret void\n"
+    "}\n"
+    "declare !callback !0 dso_local i32 @cudaLaunchKernel(i8*, i64, i32, i64, i32, i8**, i64, %struct.CUstream_st*)\n"
+    "!0 = !{!1}\n"
+    "!1 = !{i64 1, i64 0, i64 5, i1 false}\n";
+
+  std::unique_ptr<Module> M = parseIR(C, IR);
+  ASSERT_TRUE(M);
+
+  Function *Callback = M->getFunction("callback");
+  ASSERT_NE(Callback, nullptr);
+
+  const Use *CallbackUse = Callback->getSingleUndroppableUse();
+  ASSERT_NE(CallbackUse, nullptr);
+
+  AbstractCallSite ACS(CallbackUse);
+  EXPECT_TRUE(ACS);
+  EXPECT_TRUE(ACS.isCallbackCall());
+  EXPECT_TRUE(ACS.isCallee(CallbackUse));
+  EXPECT_EQ(ACS.getCalledFunction(), Callback);
+  EXPECT_TRUE(ACS.getCallArgOperand(1));
+  EXPECT_EQ(ACS.getCallArgOperand(1)->getName(), "a.addr.cast");
+}
diff --git a/llvm/unittests/IR/LegacyPassManagerTest.cpp b/llvm/unittests/IR/LegacyPassManagerTest.cpp
--- a/llvm/unittests/IR/LegacyPassManagerTest.cpp
+++ b/llvm/unittests/IR/LegacyPassManagerTest.cpp
@@ -768,7 +768,7 @@
                        "}\n"
                        "\n"
                        "!0 = !{!1}\n"
-                       "!1 = !{i64 0, i64 1, i1 false}";
+                       "!1 = !{i64 0, i64 0, i64 1, i1 false}";
 
       SMDiagnostic Err;
       std::unique_ptr<Module> M = parseAssemblyString(IR, Err, Context);