[RFC][POC] Introduce callback argument encoding mode into callback metadata

[tianshilei1992]

Summary

Current !callback metadata assumes that arguments of callback function
are encoded in broker function arguments, as the following example shows:

void callback(int, float, double);
void broker(void(*cb)(int, float, double), int, float, double);

We call the argument encoding as "flat mode". However, it is not always the
case. For some programming model, such as CUDA and OpenMP offloading, arguments
are "stacked":

void callback(int *, float *, double *);
void broker(void(*cb)(int *, float *, double *), void **args, size_t size) {
  // ...
  cb(args[0], args[1], args[2]);
}

Let's call this encode as "stacked mode" (a more approporiate name is welcomed).
Current !callback metadata cannot handle this case very well. In fact, we can
establish connection between callback function arguments and args passed to
the broker function via the use chain of args. For example, before we call the
broker function, we need to construct args by setting its every element. In
this way, we can know which pointer eventually corresponds to which callback
function argument.

In order to do that, we need to tell apart the two different parameter encoding
modes. In this patch, the !callback metadata is updated in the following way:

1. The first operand is an i64 value, representing parameter encoding mode. 0 is "flat mode", and 1 is "stacked mode".
2. The second operand is the callback function index, same as the previous first operand.
3. The next one or more operands are still parameter encoding with a little difference. If the encoding mode is 0, it is same as current way. If the mode is 1, there must only be one operand: the index of the base pointer (e.g. the index of void **args above).
4. The last operand is still for the variadic arguments.

P.S. People might argue in CUDA or OpenMP offloading, the broker function
actually doesn't accept a callback function pointer. It's usually a global
symbol which can be used to find the entry kernel function during the runtime.
More important, the kernel function (or "callback" function) cannot be seen in
the host module at all. What's the point of doing this? Yes, that's true for now.
We're working on the heterogenous module, which basically will "merge" (or "link")
kernel moduel and host module together. In this way, we can optimize them with
the information from both host and kernel side. We can bridge the kernel function
and broker function via the global symbol somehow to make the kernel function a
"virtual" callback function. Although pointers in void **args will not be used
by the kernel function directly (as shown in the example above), the real
arguments passed to the kernel function are 1:1 mapped from pointers in args.
Therefore the correspondence still exists.

[tianshilei1992]

clang/lib/Sema/SemaDeclAttr.cpp

3644

Need to revise this line to see which is the right number passed here.

llvm/include/llvm/IR/AbstractCallSite.h

55

I guess this enum should be in a more common place such that other files can use it. Any suggestion?

[tianshilei1992]

update doc in clang

[tianshilei1992]

update llvm doc

[tianshilei1992]

put the encoding mode into ParameterEncoding

[tianshilei1992]

add the support for stacked mode in AbstractCallSite::getCallArgOperand

[tianshilei1992]

We probably don't need the index for size_t size. If the ArgNo is out of range, we simply return nullptr. Besides, CUDA function cudaLaunchKernel doesn't have an argument for size as well.

[tianshilei1992]

llvm/lib/IR/AbstractCallSite.cpp

209

If we don't use have a size, GEP will be nullptr here, and we can just return it.

[tianshilei1992]

remove the requirement of size_t size

[aaron.ballman]

clang/include/clang/Basic/AttrDocs.td

432

It looks like a bunch of unrelated whitespace changes snuck in.

5221

Rather than use 0 and 1 directly, any reason not to use a named enumerator so that the user can write: __attribute__((callback(stacked, 3, 4)))?

Also, should this attribute argument be the last argument in the list and made optional so that existing code will continue to work?

clang/include/clang/Basic/DiagnosticSemaKinds.td

2940

I'm not certain it's reasonable to turn this into an error given that this attribute already exists in the wild. This basically breaks all users of the attribute.

clang/include/clang/Basic/AttrDocs.td

 that appears to be capable of returning to its caller.
   }];
 }
 
 def NoMergeDocs : Documentation {
   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.
   }];
 }
 
 def AssertCapabilityDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "assert_capability, assert_shared_capability";
   let Content = [{

 
 
 def NullabilityDocs : DocumentationCategory<"Nullability Attributes"> {
   let Content = [{
 Whether a particular pointer may be "null" is an important concern when working
 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
 be null (``_Nonnull``).
 
 The nullability (type) qualifiers express whether a value of a given pointer
 type can be null (the ``_Nullable`` qualifier), doesn't have a defined meaning
 for null (the ``_Nonnull`` qualifier), or for which the purpose of null is
memory is not available rather than returning a null pointer:
 
   .. code-block:: c
 
     extern void * malloc (size_t size) __attribute__((returns_nonnull));
 
 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
 }];
 }
 
 def NoAliasDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 The ``noalias`` attribute indicates that the only memory accesses inside

 }
 
 def CallbackDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 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.
 
 The ``callback`` attribute, which is directly translated to ``callback``
 metadata <http://llvm.org/docs/LangRef.html#callback-metadata>, make the
 connection between the call to the annotated function and the callback callee.
 This can enable interprocedural optimizations which were otherwise impossible.
 If a function parameter is mentioned in the ``callback`` attribute, through its
 position, it is undefined if that parameter is used for anything other than the
 actual callback. Inspected, captured, or modified parameters shall not be
 listed in the ``callback`` metadata.
 
 Example encodings for the callback performed by ``pthread_create`` are shown
 below. The explicit attribute annotation indicates that the third parameter
 (``start_routine``) is called zero or more times by the ``pthread_create`` function,
 and that the fourth parameter (``arg``) is passed along. Note that the callback
 behavior of ``pthread_create`` is automatically recognized by Clang. In addition,
 the declarations of ``__kmpc_fork_teams`` and ``__kmpc_fork_call``, generated for
 ``#pragma omp target teams`` and ``#pragma omp parallel``, respectively, are also
 automatically recognized as broker functions. Further functions might be added
 in the future.
 
   .. 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);
+
   }];
 }
 
 def CalledOnceDocs : Documentation {
   let Category = DocCatVariable;
   let Content = [{
 The ``called_once`` attribute specifies that the annotated function or method
 parameter is invoked exactly once on all execution paths. It only applies

 }];
 }
 
 def CFGuardDocs : Documentation {
   let Category = DocCatFunction;
   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>`
 }];
 }
 
 def CUDADeviceBuiltinSurfaceTypeDocs : Documentation {
   let Category = DocCatType;

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_format_attribute_requires_variadic : Error<
   "format attribute requires variadic function">;
 def err_format_attribute_not : Error<"format argument not %0">;
 def err_format_attribute_result_not : Error<"function does not return %0">;
 def err_format_attribute_implicit_this_format_string : Error<
   "format attribute cannot specify the implicit this argument as the format "
   "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<
   "multiple 'callback' attributes specified">;
 def err_callback_attribute_argument_unknown : Error<
   "'callback' attribute argument %0 is not a known function parameter">;
 def err_callback_callee_no_function_type : Error<
   "'callback' attribute callee does not have function type">;

clang/lib/CodeGen/CodeGenModule.cpp

if (!CodeGenOpts.SanitizeCfiCrossDso ||
       !CodeGenOpts.SanitizeCfiCanonicalJumpTables)
     CreateFunctionTypeMetadataForIcall(FD, F);
 
   if (getLangOpts().OpenMP && FD->hasAttr<OMPDeclareSimdDeclAttr>())
     getOpenMPRuntime().emitDeclareSimdFunction(FD, F);
 
   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,
+                                    EncodingMode, CalleeIdx, PayloadIndices,
   }
 }
 
 void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
   assert((isa<llvm::Function>(GV) || !GV->isDeclaration()) &&
          "Only globals with definition can force usage.");
   LLVMUsed.emplace_back(GV);
 }

clang/lib/Sema/SemaDeclAttr.cpp

static void handleFormatAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   FormatAttr *NewAttr = S.mergeFormatAttr(D, AL, II, Idx, FirstArg);
   if (NewAttr)
     D->addAttr(NewAttr);
 }
 
 /// 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;
   }
 
   bool HasImplicitThisParam = isInstanceMethod(D);
   int32_t NumArgs = getFunctionOrMethodNumParams(D);
 
if (AL.isArgIdent(I)) {
         return;
       }
 
       SR = IdxExpr->getSourceRange();
     } else {
       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;
     }
 
     // Adjust for the case we do not have an implicit "this" parameter. In this
     // case we decrease all positive values by 1 to get LLVM argument indices.
     if (!HasImplicitThisParam && ArgIdx > 0)
       ArgIdx -= 1;
 
     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.
   if (CalleeIdx < (int)HasImplicitThisParam) {
     S.Diag(AL.getLoc(), diag::err_callback_attribute_invalid_callee)
         << AL.getRange();
     return;
   }
static void handleCallbackAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   const auto *CalleeFnProtoType = dyn_cast<FunctionProtoType>(CalleeFnType);
   if (!CalleeFnProtoType) {
     S.Diag(AL.getLoc(), diag::err_callback_callee_no_function_type)
         << AL.getRange();
     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;
   }
 
   if (CalleeFnProtoType->isVariadic()) {
     S.Diag(AL.getLoc(), diag::err_callback_callee_is_variadic) << AL.getRange();
     return;
   }
 

clang/test/CodeGen/attr-callback.c

 // RUN: %clang_cc1 -triple i386-unknown-unknown %s -emit-llvm -o - | FileCheck %s
 
 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}

clang/test/CodeGen/callback_annotated.c

 // 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;
 }
 
 static int ThreeInt2Int(int a, int b, int c) {
   return a * b + c;
 }
 
 void foo() {
   broker0(VoidPtr2VoidPtr, 0l);
   broker1(0l, VoidPtr2VoidPtr);
   broker2(foo);
   broker3(1, 4, 5, ThreeInt2Int, 1);
   broker4(4, 2, 7, ThreeInt2Int, 0);
   broker5(8, 0, 3, ThreeInt2Int, 4);
 }
 
 // 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}

clang/test/CodeGen/callback_openmp.c

 // RUN: %clang_cc1 -triple i386-unknown-unknown -fopenmp %s -emit-llvm -o - -disable-llvm-optzns | FileCheck %s
 
 // 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);
 void work12(int, int);
 
 void foo(int q) {
   int p = 2;
 

clang/test/CodeGenCXX/attr-callback.cpp

 // RUN: %clang_cc1 -triple i386-unknown-unknown %s -emit-llvm -o - | FileCheck %s
 
 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]+]]
 void Base::this_unknown_this(void (*callback)(Base *, Base *, Base *)) {
 }
 
 struct Derived_1 : public Base {
-  __attribute__((callback(1))) virtual void
+  __attribute__((callback(0, 1))) virtual void
   virtual_1(void (*callback)(void)) override;
 };
 
 // CHECK-DAG:      define{{.*}} void @_ZN9Derived_19virtual_1EPFvvE({{[^!]*!callback}} ![[cid0]]
 void Derived_1::virtual_1(void (*callback)(void)) {}
 
 struct Derived_2 : public Base {
   void virtual_1(void (*callback)(void)) override;
 };
 
 // CHECK-DAG: define{{.*}} void @_ZN9Derived_29virtual_1EPFvvE
 // CHECK-NOT: !callback
 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}

clang/test/OpenMP/parallel_codegen.cpp

 // ALL: define linkonce_odr {{.*}}void [[FOO1]](i8** %argc)
 // CHECK-DEBUG-DAG: define linkonce_odr void [[FOO1]](i8** %argc)
 // CHECK-DEBUG-DAG: define internal void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i8*** nonnull align {{[0-9]+}} dereferenceable({{4|8}}) %argc, i64 %{{.+}})
 // CHECK-DEBUG-DAG: call void [[OMP_OUTLINED_DEBUG]]({{[^)]+}}){{[^,]*}}, !dbg
 
 // 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

clang/test/Sema/attr-callback-broken.c

 // 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(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(-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, 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(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(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}}

clang/test/Sema/attr-callback.c

 // RUN: %clang_cc1 %s -verify -fsyntax-only
 
 // 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);

clang/test/SemaCXX/attr-callback.cpp

 // RUN: %clang_cc1 %s -verify -fsyntax-only
 
 // expected-no-diagnostics
 
 class C_in_class {
 #include "../Sema/attr-callback.c"
 };
 
 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)) {
 }
 
 void Base::no_args_2(void (*callback)(void)) {
 }
 
 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
   virtual_3(void (*callback)(void)) override {}
 };
 
 struct Derived_2 : public Base {
 
-  __attribute__((callback(callback))) virtual void
+  __attribute__((callback(0, callback))) virtual void
   virtual_1(void (*callback)(void)) override;
 
   virtual void
   virtual_2(void (*callback)(void)) override;
 
   virtual void
   virtual_3(void (*callback)(void)) override;
 };
 
 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)) {}

llvm/docs/LangRef.rst

 be passed to the specified callback function still hold, even if the callback
 function is not dynamically invoked. The broker is allowed to invoke the
 callback function more than once per invocation of the broker. The broker is
 also allowed to invoke (directly or indirectly) the function passed as a
 callback through another use. Finally, the broker is also allowed to relay the
 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
    have misc.highlighting_failure set?
 
 .. code-block:: text
 
     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
    have misc.highlighting_failure set?
 
 .. code-block:: text
 
     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}
 
 
 '``unpredictable``' Metadata
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
 ``unpredictable`` metadata may be attached to any branch or switch
 instruction. It can be used to express the unpredictability of control

 
 Syntax:
 """""""
 
 ::
 
       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)
 
 Overview:
 """""""""
 
 The statepoint intrinsic represents a call which is parse-able by the

 implementation limitation.  Other than the type of the return value,
 the semantics are the same.
 
 Operands:
 """""""""
 
 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:
 """"""""""
 
 The ``gc.result`` represents the return value of the call target of
 the ``statepoint``.  The type of the ``gc.result`` must exactly match
 the type of the target.  If the call target returns void, there will
 be no ``gc.result``.
 
 A ``gc.result`` is modeled as a 'readnone' pure function.  It has no
 side effects since it is just a projection of the return value of the
 previous call represented by the ``gc.statepoint``.
 
 'llvm.experimental.gc.relocate' Intrinsic
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
 Syntax:
 """""""
 
 ::
 
       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:
 """""""""
 
 A ``gc.relocate`` returns the potentially relocated value of a pointer
 at the safepoint.
 
 Operands:
 """""""""
 
 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
 corresponding statepoint's :ref:`gc-live <ob_gc_live>` operand bundle.
 
 The second argument is an index which specifies the allocation for the pointer
 being relocated. The associated value must be within the object with which the
 pointer being relocated is associated. The optimizer is free to change *which*

llvm/include/llvm/IR/AbstractCallSite.h

 /// to note that the mapping is often partial. Thus, some arguments of the
 /// call/invoke instruction are mapped to parameters of the callee while others
 /// are not.
 class AbstractCallSite {
 public:
 
   /// 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
     /// element of the encoding vector represents which argument of the call
     /// site CB is the callback callee. The remaining elements map parameters
     /// (identified by their position) to the arguments that will be passed
     /// through (also identified by position but in the call site instruction).
     ///
     /// NOTE that we use LLVM argument numbers (starting at 0) and not
     /// clang/source argument numbers (starting at 1). The -1 entries represent
     /// unknown values that are passed to the callee.
     using ParameterEncodingTy = SmallVector<int, 0>;
     ParameterEncodingTy ParameterEncoding;
-
   };
 
 private:
 
   /// The underlying call site:
   ///   caller -> callee,             if this is a direct or indirect call site
   ///   caller -> broker function,    if this is a callback call site
   CallBase *CB;
assert(!CI.ParameterEncoding.empty() &&
            "Callback without parameter encoding!");
 
     // If the use is actually in a constant cast expression which itself
     // has only one use, we look through the constant cast expression.
     if (auto *CE = dyn_cast<ConstantExpr>(U->getUser()))
       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
   /// Arg.
   int getCallArgOperandNo(Argument &Arg) const {
     return getCallArgOperandNo(Arg.getArgNo());
   }
 
   /// Return the operand index of the underlying instruction associated with
   /// the function parameter number @p ArgNo or -1 if there is none.
   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.
   Value *getCallArgOperand(Argument &Arg) const {
     return getCallArgOperand(Arg.getArgNo());
   }
 
   /// 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
   /// valid for callback calls!
   const Use &getCalleeUseForCallback() const {
     int CalleeArgIdx = getCallArgOperandNoForCallee();
     assert(CalleeArgIdx >= 0 &&
            unsigned(CalleeArgIdx) < getInstruction()->getNumOperands());

llvm/include/llvm/IR/MDBuilder.h

public:
   /// calls.
   MDNode *createCallees(ArrayRef<Function *> Callees);
 
   //===------------------------------------------------------------------===//
   // Callback metadata.
   //===------------------------------------------------------------------===//
 
   /// 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.
   MDNode *mergeCallbackEncodings(MDNode *ExistingCallbacks, MDNode *NewCB);
 
   //===------------------------------------------------------------------===//
   // AA metadata.
   //===------------------------------------------------------------------===//

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

if (FnID == OMPRTL___kmpc_fork_call || FnID == OMPRTL___kmpc_fork_teams) {
       if (!Fn->hasMetadata(LLVMContext::MD_callback)) {
         LLVMContext &Ctx = Fn->getContext();
         MDBuilder MDB(Ctx);
         // Annotate the callback behavior of the runtime function:
         //  - The callback callee is argument number 2 (microtask).
         //  - 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,
+        Fn->addMetadata(LLVMContext::MD_callback,
                         *MDNode::get(Ctx, {MDB.createCallbackEncoding(
-                                  2, {-1, -1}, /* VarArgsArePassed */ true)}));
+                                              0, 2, {-1, -1},
+                                              /* VarArgsArePassed */ true)}));
       }
     }
 
     LLVM_DEBUG(dbgs() << "Created OpenMP runtime function " << Fn->getName()
                       << " with type " << *Fn->getFunctionType() << "\n");
     addAttributes(FnID, *Fn);
 
   } else {
if (!F->hasMetadata(llvm::LLVMContext::MD_callback)) {
       // Annotate the callback behavior of the __kmpc_fork_call:
       //  - The callback callee is argument number 2 (microtask).
       //  - The first two arguments of the callback callee are unknown (-1).
       //  - All variadic arguments to the __kmpc_fork_call are passed to the
       //    callback callee.
       F->addMetadata(
           llvm::LLVMContext::MD_callback,
           *llvm::MDNode::get(
-              Ctx, {MDB.createCallbackEncoding(2, {-1, -1},
+              Ctx, {MDB.createCallbackEncoding(0, 2, {-1, -1},
                                                /* VarArgsArePassed */ true)}));
     }
   }
 
   OutlineInfo OI;
   OI.PostOutlineCB = [=](Function &OutlinedFn) {
     // Add some known attributes.
     OutlinedFn.addParamAttr(0, Attribute::NoAlias);

llvm/lib/IR/AbstractCallSite.cpp

 //
 // For more information see:
 // https://llvm.org/devmtg/2018-10/talk-abstracts.html#talk20
 //
 //===----------------------------------------------------------------------===//
 
 #include "llvm/IR/AbstractCallSite.h"
 #include "llvm/ADT/Statistic.h"
+#include "llvm/IR/Instructions.h"
 #include "llvm/Support/Debug.h"
 
 using namespace llvm;
 
 #define DEBUG_TYPE "abstract-call-sites"
 
 STATISTIC(NumCallbackCallSites, "Number of callback call sites created");
 STATISTIC(NumDirectAbstractCallSites,
if (!Callee)
     return;
 
   MDNode *CallbackMD = Callee->getMetadata(LLVMContext::MD_callback);
   if (!CallbackMD)
     return;
 
   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())
       CallbackUses.push_back(CB.arg_begin() + CBCalleeIdx);
   }
 }
 
 /// Create an abstract call site from a use.
if (!CallbackMD) {
     CB = nullptr;
     return;
   }
 
   unsigned UseIdx = CB->getArgOperandNo(U);
   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)
       continue;
     CallbackEncMD = OpMD;
     break;
   }
 
   if (!CallbackEncMD) {
     NumInvalidAbstractCallSitesNoCallback++;
     CB = nullptr;
     return;
   }
 
   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);
     assert(OpAsCM->getType()->isIntegerTy(64) &&
            "Malformed !callback metadata");
 
     int64_t Idx = cast<ConstantInt>(OpAsCM->getValue())->getSExtValue();
     assert(-1 <= Idx && Idx <= NumCallOperands &&
AbstractCallSite::AbstractCallSite(const Use *U)
 
   if (VarArgFlagAsCM->getValue()->isNullValue())
     return;
 
   // Add all variadic arguments at the end.
   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);
+  }
+}

llvm/lib/IR/MDBuilder.cpp

 
 MDNode *MDBuilder::createCallees(ArrayRef<Function *> Callees) {
   SmallVector<Metadata *, 4> Ops;
   for (Function *F : Callees)
     Ops.push_back(createConstant(F));
   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)
     Ops.push_back(createConstant(ConstantInt::get(Int64, ArgNo, true)));
 
   Type *Int1 = Type::getInt1Ty(Context);
   Ops.push_back(createConstant(ConstantInt::get(Int1, VarArgArePassed)));
 

llvm/test/Analysis/CallGraph/callback-calls.ll

 
 define void @callback(i32* %arg) {
   ret void
 }
 
 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}

llvm/test/Analysis/CallGraph/ignore-callback-uses.ll

 omp.par.outlined.exit:                            ; preds = %omp.par.pre_finalize
   br label %omp.par.exit.split.exitStub
 }
 
 ; Function Attrs: nounwind
 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}

llvm/test/Transforms/Attributor/IPConstantProp/multiple_callbacks.ll

entry:
   call void @broker(i32 (i32)* nonnull @cb1, i32 (i32)* nonnull @cb2, i32 (i32)* nonnull @cb2, i32 0, i32 1)
   call void @broker(i32 (i32)* nonnull @cb3, i32 (i32)* nonnull @cb2, i32 (i32)* nonnull @cb3, i32 0, i32 1)
   call void @broker(i32 (i32)* nonnull @cb4, i32 (i32)* nonnull @cb4, i32 (i32)* nonnull @cb4, i32 0, i32 1)
   ret void
 }
 
 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 }
 ;.
 ; 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}
 ;.

llvm/test/Transforms/Attributor/IPConstantProp/openmp_parallel_for.ll

 declare dso_local void @__kmpc_for_static_init_4(%struct.ident_t*, i32, i32, i32*, i32*, i32*, i32*, i32, i32)
 
 declare dso_local void @bar(i32, float, double)
 
 declare dso_local void @__kmpc_for_static_fini(%struct.ident_t*, i32)
 
 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}
 ;.

llvm/test/Transforms/Attributor/IPConstantProp/pthreads.ll

 ; IS__CGSCC____-SAME: (i8* noalias nofree noundef nonnull readnone returned align 8 dereferenceable(1) "no-capture-maybe-returned" [[ARG:%.*]]) #[[ATTR0]] {
 ; IS__CGSCC____-NEXT:  entry:
 ; IS__CGSCC____-NEXT:    ret i8* [[ARG]]
 ;
 entry:
   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 }
 ;.
 ; 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}
 ;.

llvm/test/Transforms/Attributor/IPConstantProp/thread_local_acs.ll

 ;
 entry:
   call void @broker(i32* nonnull @gtl, i32 (i32*, i32*)* nonnull @callee, i32* nonnull @gsh)
   ret void
 }
 
 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 }
 ;.
 ; 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}
 ;.

llvm/test/Transforms/Attributor/callbacks.ll

entry:
   ret void
 }
 
 declare void @t3_check(i32* nocapture align 256, i64, i32* nocapture)
 
 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}
 ;.

llvm/test/Transforms/Attributor/noundef.ll

 ; CHECK-NEXT:    call void (i1, ...) @callee_with_dead_arg(i1 undef)
 ; CHECK-NEXT:    ret void
 ;
   call void (i1, ...) @callee_with_dead_arg(i1 noundef true)
   ret void
 }
 
 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}
 ;.

llvm/test/Transforms/OpenMP/parallel_deletion.ll

 
 declare void @readonly() readonly
 
 declare void @readnone() readnone
 
 !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}
 !7 = !{!"Simple C/C++ TBAA"}

llvm/test/Transforms/OpenMP/parallel_deletion_cg_update.ll

 ; RUN: opt < %s -instcombine -attributor-cgscc -print-callgraph -disable-output -verify 2>&1 -enable-new-pm=0 | FileCheck %s
 
 ; CHECK: Call graph node <<null function>><<{{.*}}>>  #uses=0
 ; CHECK:   CS<None> calls function 'dead_fork_call'
 ; CHECK:   CS<None> calls function '.omp_outlined..0'
 ; CHECK:   CS<None> calls function '__kmpc_fork_call'
 ; CHECK:   CS<None> calls function 'live_fork_call'
 ; CHECK:   CS<None> calls function '.omp_outlined..1'
 ; 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'
 ;
 ; CHECK: Call graph node for function: '__kmpc_fork_call'<<{{.*}}>>  #uses=3
 ; CHECK:   CS<None> calls external node
 ;
 ; CHECK: Call graph node for function: 'd'<<{{.*}}>>  #uses=2
 ; CHECK:   CS<None> calls external node
entry:
   ret void
 }
 
 declare dso_local void @d(...)
 
 !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}

llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll

 !16 = !DISubroutineType(types: !17)
 !17 = !{null}
 !18 = !DILocation(line: 8, column: 1, scope: !15)
 !19 = !DILocation(line: 10, column: 1, scope: !15)
 !20 = !DILocation(line: 12, column: 1, scope: !15)
 !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}
 !28 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !29)
 !29 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !30)
 !30 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !31, size: 64)
 !31 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !32)
 !32 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)

llvm/test/Transforms/OpenMP/parallel_region_merging.ll

entry:
   ret void
 }
 
 
 !llvm.module.flags = !{!0}
 
 !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:
 ; CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
 ; CHECK-NEXT:    store i32 [[A]], i32* [[A_ADDR]], align 4
 ; CHECK-NEXT:    [[OMP_GLOBAL_THREAD_NUM:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* [[GLOB1:@.*]])
 ; CHECK-NEXT:    br label [[OMP_PARALLEL:%.*]]
 ; CHECK:       omp_parallel:

llvm/unittests/IR/AbstractCallSiteTest.cpp

 static std::unique_ptr<Module> parseIR(LLVMContext &C, const char *IR) {
   SMDiagnostic Err;
   std::unique_ptr<Module> Mod = parseAssemblyString(IR, Err, C);
   if (!Mod)
     Err.print("AbstractCallSiteTests", errs());
   return Mod;
 }
 
-TEST(AbstractCallSite, CallbackCall) {
+TEST(AbstractCallSite, FlatCallbackCall) {
   LLVMContext C;
 
   const char *IR =
       "define void @callback(i8* %X, i32* %A) {\n"
       "  ret void\n"
       "}\n"
       "define void @foo(i32* %A) {\n"
       "  call void (i32, void (i8*, ...)*, ...) @broker(i32 1, void (i8*, ...)* bitcast (void (i8*, i32*)* @callback to void (i8*, ...)*), i32* %A)\n"
       "  ret void\n"
       "}\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);
 
   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);
 }
+
+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");
+}

llvm/unittests/IR/LegacyPassManagerTest.cpp

TEST(PassManager, CallbackCallsModifier0) {
                        "  ret void\n"
                        "}\n"
                        "\n"
                        "define internal void @callback1(i8* %arg) {\n"
                        "  ret void\n"
                        "}\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);
       if (!M)
         Err.print("LegacyPassManagerTest", errs());
 
       CallbackCallsModifierPass *P = new CallbackCallsModifierPass();
       legacy::PassManager Passes;