[Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values

[skc7]

Summary

Add the ability to put attribute((maybe_undef)) on function arguments. Clang will now introduce a freeze instruction on the argument. Previous review relating to the noundef attribute issue: D128907.

llvm discourse RFC topic: RFC

[aaron.ballman]

I'm in C standards meetings this week and don't have a lot of ability to thoroughly review this yet, but the most immediate concern which springs to mind for me is that this is exposing LLVM implementation details to users. Users should not have to think about things in terms of LLVM IR markings like poison and undef, and I worry that this is an expert-only feature that will be easy to misuse and introduce security issues. It sounds like there's an extremely specific use case in mind for this attribute; is there a reason why this needs to be the user's problem instead of something the compiler can infer or handle on the backend?

[arsenm]

In D130224#3668225, @aaron.ballman wrote:

It sounds like there's an extremely specific use case in mind for this attribute; is there a reason why this needs to be the user's problem instead of something the compiler can infer or handle on the backend?

The intended user is a handful of header wrapper functions around intrinsics, not general users

[aaron.ballman]

In D130224#3668240, @arsenm wrote:

In D130224#3668225, @aaron.ballman wrote:

It sounds like there's an extremely specific use case in mind for this attribute; is there a reason why this needs to be the user's problem instead of something the compiler can infer or handle on the backend?

The intended user is a handful of header wrapper functions around intrinsics, not general users

Great, then it sounds like we don't need to expose this as an attribute because we know which intrinsics need special behavior, so we can handle them specially at codegen time or within the backend?

[arsenm]

In D130224#3668243, @aaron.ballman wrote:

In D130224#3668240, @arsenm wrote:

In D130224#3668225, @aaron.ballman wrote:

It sounds like there's an extremely specific use case in mind for this attribute; is there a reason why this needs to be the user's problem instead of something the compiler can infer or handle on the backend?

The intended user is a handful of header wrapper functions around intrinsics, not general users

Great, then it sounds like we don't need to expose this as an attribute because we know which intrinsics need special behavior, so we can handle them specially at codegen time or within the backend?

The problem is the wrappers. We originally talked about specifically handling the builtins, but that falls apart given the real users come from wrapper functions which look like ordinary code

[skc7]

Rebase. Remove skipping noundef attribute based on maybe_undef.

[nhaehnle]

In D130224#3668225, @aaron.ballman wrote:

I'm in C standards meetings this week and don't have a lot of ability to thoroughly review this yet, but the most immediate concern which springs to mind for me is that this is exposing LLVM implementation details to users. Users should not have to think about things in terms of LLVM IR markings like poison and undef, and I worry that this is an expert-only feature that will be easy to misuse and introduce security issues.

Here's how I would tentatively describe the attribute in terms that mesh better with how I understand C and C++:

As an exception to the rule that loading from an unitialized variable is undefined behavior, if the loaded value is used immediately as an __attribute__((maybe_undef)) argument in a function call, the loaded value is implementation-defined. It may vary between multiple runs of the program, and it may vary between multiple uses of the uninitialized variable.

This requires no thinking about LLVM IR and undef/poison.

There may have to be some caveats about bools and enums (generally, types where not all possible values in memory are actually legal values of the type), but I don't know enough about those language standards to judge that.

clang/include/clang/Basic/AttrDocs.td

276

param shouldn't be of void type, right?

[jdoerfert]

Generating freeze seems an option, just not sure we can easily get all locations with this approach.
Anyway, some code comments below.

clang/lib/CodeGen/CGCall.cpp

2061

Early exists please.

if (!TargetDecl)

return false

no ArgHasMayBeUndefAttr is needed, just return true/false, etc.

4906

Style: Val, also below.

clang/test/CodeGenHIP/maybe_undef-attr-verify.hip

34

Note, not necessarily to address: I would have preferred auto generated check lines, should not be many and it's easier to see what's going on, IMHO.

[skc7]

Fixes as per review comments.

[skc7]

Rebase. Ping.

[aaron.ballman]

In D130224#3668267, @arsenm wrote:

In D130224#3668243, @aaron.ballman wrote:

In D130224#3668240, @arsenm wrote:

In D130224#3668225, @aaron.ballman wrote:

It sounds like there's an extremely specific use case in mind for this attribute; is there a reason why this needs to be the user's problem instead of something the compiler can infer or handle on the backend?

The intended user is a handful of header wrapper functions around intrinsics, not general users

Great, then it sounds like we don't need to expose this as an attribute because we know which intrinsics need special behavior, so we can handle them specially at codegen time or within the backend?

The problem is the wrappers. We originally talked about specifically handling the builtins, but that falls apart given the real users come from wrapper functions which look like ordinary code

I'm still not seeing the issue fully. My understanding of the situation (which may be wrong) is that Clang lowers to LLVM IR and adds noundef markings at call sites that this patch attempts to let the user then undo. However, why can Clang's CodeGen not notice the special builtin and walk up the call chain to mark all the wrapper functions to ensure everything is marked appropriately? There might be a small perf loss (if there are wrappers around wrappers around wrappers kind of situation), but this means we don't have to expose this problem to users.

In D130224#3672338, @nhaehnle wrote:

In D130224#3668225, @aaron.ballman wrote:

I'm in C standards meetings this week and don't have a lot of ability to thoroughly review this yet, but the most immediate concern which springs to mind for me is that this is exposing LLVM implementation details to users. Users should not have to think about things in terms of LLVM IR markings like poison and undef, and I worry that this is an expert-only feature that will be easy to misuse and introduce security issues.

Here's how I would tentatively describe the attribute in terms that mesh better with how I understand C and C++:

As an exception to the rule that loading from an unitialized variable is undefined behavior, if the loaded value is used immediately as an __attribute__((maybe_undef)) argument in a function call, the loaded value is implementation-defined. It may vary between multiple runs of the program, and it may vary between multiple uses of the uninitialized variable.

This requires no thinking about LLVM IR and undef/poison.

Thanks, this helps somewhat (though I would argue that the loaded value is *not* implementation-defined but is instead unspecified, because that matches the semantics you mention about not needing to pick the same loaded value every time). However, this design is basically asking the user to decide when something should or should not be undefined behavior, which means this attribute will almost certainly be abused under the guise of "this makes my program safer because it removes UB". I still think that a solution which doesn't require the user to make such dangerous decisions on their own is a better design if we can accomplish it.

[arsenm]

In D130224#3668243, @aaron.ballman wrote:

I'm still not seeing the issue fully. My understanding of the situation (which may be wrong) is that Clang lowers to LLVM IR and adds noundef markings at call sites that this patch attempts to let the user then undo. However, why can Clang's CodeGen not notice the special builtin and walk up the call chain to mark all the wrapper functions to ensure everything is marked appropriately? There might be a small perf loss (if there are wrappers around wrappers around wrappers kind of situation), but this means we don't have to expose this problem to users.

This requires maintaining logic in clang to specially recognize these N wrapper functions, and there would be no indication in the source that this magic is occurring. It's less maintainable because it requires tight, magic coupling between these functions in the source and special handling in clang to be maintained if anything were to change

[jdoerfert]

I'm still not seeing the issue fully. My understanding of the situation (which may be wrong) is that Clang lowers to LLVM IR and adds noundef markings at call sites that this patch attempts to let the user then undo. However, why can Clang's CodeGen not notice the special builtin and walk up the call chain to mark all the wrapper functions to ensure everything is marked appropriately? There might be a small perf loss (if there are wrappers around wrappers around wrappers kind of situation), but this means we don't have to expose this problem to users.

I think you are right, we could identify shuffles and go up the call chain in clang (codegen) if we wanted.

I still think that a solution which doesn't require the user to make such dangerous decisions on their own is a better design if we can accomplish it.

This is a valid opinion but I don't see how this matches what we do (or has been done) in the past. I mean, there are plenty of command line options and attributes that do exactly this, let the user "define", or better "refine", semantics.
The entire "initialize with 0" stuff comes to mind, making new things UB with -ffast-math and friends, overloadable, etc.

---

All that said, I'm generally more in favor of explicit annotations in the source than implicit compiler magic to implement a special case.
Others should chime in.

[aaron.ballman]

In D130224#3677144, @arsenm wrote:

In D130224#3668243, @aaron.ballman wrote:

I'm still not seeing the issue fully. My understanding of the situation (which may be wrong) is that Clang lowers to LLVM IR and adds noundef markings at call sites that this patch attempts to let the user then undo. However, why can Clang's CodeGen not notice the special builtin and walk up the call chain to mark all the wrapper functions to ensure everything is marked appropriately? There might be a small perf loss (if there are wrappers around wrappers around wrappers kind of situation), but this means we don't have to expose this problem to users.

This requires maintaining logic in clang to specially recognize these N wrapper functions, and there would be no indication in the source that this magic is occurring. It's less maintainable because it requires tight, magic coupling between these functions in the source and special handling in clang to be maintained if anything were to change

I was thinking of something more general than that where Clang only has to specially recognize the builtins and which arguments to the builtin should be tracked to make maybeundef when walking up the call chain. However, you're not wrong about this being a bit magical in terms of behavior. I don't know that we have anything else which behaves this way.

In D130224#3677167, @jdoerfert wrote:

I still think that a solution which doesn't require the user to make such dangerous decisions on their own is a better design if we can accomplish it.

This is a valid opinion but I don't see how this matches what we do (or has been done) in the past. I mean, there are plenty of command line options and attributes that do exactly this, let the user "define", or better "refine", semantics.
The entire "initialize with 0" stuff comes to mind, making new things UB with -ffast-math and friends, overloadable, etc.

Thanks, this is a good question to be asking!

We are adding attributes in a way which doesn't scale well; many attributes interact with one another and we have basically no test coverage for those situations and very little thought put into diagnostic behavior. Now that we've got over 400 attributes in the tree, I'm realizing I probably should have been saying "no" more to attributes that don't carry enough weight to be worth the maintenance burden, but oh well, I can't unring that bell. I have *no issue* with attributes that are going to be things we expect users to use themselves. Stuff like "initialize with 0", overloadable, etc are all things that enable users to do something they couldn't previously do. However, in recent history, we've been adding more attributes that are for compiler glue code rather than users. On the one hand, attributes are a great use for such a thing. On the other hand, there's no way to control who uses the attribute (unless we put in special logic to only allow it to be written in a system header, or something like that). So glue code attributes are a different "thing" to me than user-facing attributes, and I'm worried at how many we're adding without any real design thought put behind them other than "we need it to solve this one case". If we could come up with some tablegen marking to say "this attribute is for glue code use only" that we could use to generate diagnostics on "misuse", generate documentation differently from, etc, that would go a long ways towards alleviating the concerns about user confusion with glue code attributes and "real" attributes (for lack of a better term).

Of course, these are high-level worries and not necessarily specific to this review or anything I expect someone here to take action on. Take this as background information rather than a code owner mandate or anything like that. Specific to this review, my concern is that this is an attribute that controls undefined behavior in LLVM rather directly (as opposed to things like "initialize to 0" which impact UB, but aren't as tightly tied to the LLVM backend architecture). That's a worry for two reasons: the tight coupling to LLVM IR attribute semantics (basically, this means LLVM's "freeze" instruction can't change meaning because it will potentially break user code) and misuse due to misunderstanding the purpose of the attribute ("this makes my code more secure!"). Special compiler magic alleviates both of those concerns because it means we aren't exposing the tight coupling and users aren't given a chance to misuse the attribute.

However, what I think I'm hearing from this thread is that there are alternative approaches that have been thought about but not tried, we're not certain how feasible those approaches are in practice, but we expect them to be materially worse than what's proposed here. So it's not "this was the path of least resistance", but "this is the correct design." Do others agree with that assessment?

[arsenm]

In D130224#3677224, @aaron.ballman wrote:

However, what I think I'm hearing from this thread is that there are alternative approaches that have been thought about but not tried, we're not certain how feasible those approaches are in practice, but we expect them to be materially worse than what's proposed here. So it's not "this was the path of least resistance", but "this is the correct design." Do others agree with that assessment?

I think this is just the least bad on the menu of available options. I don't like it, but it at least provides documentation about this special behavior.

I think the only other plausible option is to assert this is still undefined behavior and force users to update their (newly declared invalid) code. We could at least partially re-optimize to uninitialized values in the backend (although this is apparently difficult in some situations)

[aaron.ballman]

In D130224#3677261, @arsenm wrote:

In D130224#3677224, @aaron.ballman wrote:

However, what I think I'm hearing from this thread is that there are alternative approaches that have been thought about but not tried, we're not certain how feasible those approaches are in practice, but we expect them to be materially worse than what's proposed here. So it's not "this was the path of least resistance", but "this is the correct design." Do others agree with that assessment?

I think this is just the least bad on the menu of available options. I don't like it, but it at least provides documentation about this special behavior.

I think the only other plausible option is to assert this is still undefined behavior and force users to update their (newly declared invalid) code. We could at least partially re-optimize to uninitialized values in the backend (although this is apparently difficult in some situations)

First off, thank you for the great discussion about this both on and off lists while I wrapped my head around the design space and challenges involved. I appreciate your willingness to explore alternatives with me!

I think I agree that this is the least bad option on the menu. The code the user wrote is undefined behavior, and the CUDA/HIP standards would like to pretend otherwise in some circumstances. Given that users will have code invalidated if we stick to the hard line language rules, I think it's better than we remove the UB than punish the user and hope they can write a more convoluted but correct form.

In terms of the patch itself, I've left some comments. But this is also missing all semantic tests (that the attribute is diagnosed when written on something other than a parameter, accept no arguments, etc). It also is missing tests for some interesting use cases:

void func(int param);

void other() {
  int derp;
  func(derp); // How should the call behave?
}

void func(__attribute__((maybe_undef)) int param) { ... }

// Also:
void foo(__attribute__((maybe_undef)) int param);
void foo(int param) { ... } // Verify this parameter behaves as expected

void bar() {
  int derp;
  foo(derp);
}

It'd probably make sense to have a case involving templates as well to ensure that instantiation properly retains the parameter attribute, but we might already have test coverage for that (parameter attributes are pretty uncommon, but not unheard of).

clang/include/clang/Basic/Attr.td

2026

clang/include/clang/Basic/AttrDocs.td

260

261

I think this should be DocCatVariable temporarily, but we probably should add a new DocCatParameter category at some point and use that instead (in case you're feeling ambitious).

263

265

clang/lib/CodeGen/CGCall.cpp

2050–2066

One question I have is whether you ever need to mark the variadic arguments as being maybe undef. e.g., void func(int i, ...); do you need to signal that arguments passed to ... are maybe undef?

[skc7]

Changes for code review comments from @aaron.ballman

[skc7]

clang/lib/CodeGen/CGCall.cpp

2050–2066

Current change assumes variadic arguments will not have "maybe_undef" attribute. If its a function attribute, variadic arguments can inherit them (Have seen such cases in clang codebase). But "maybe_undef" is function argument attribute and I'm not sure on how to add it to variadic arguments.

[aaron.ballman]

clang/lib/CodeGen/CGCall.cpp

2050–2066

Current change assumes variadic arguments will not have "maybe_undef" attribute. If its a function attribute, variadic arguments can inherit them (Have seen such cases in clang codebase). But "maybe_undef" is function argument attribute and I'm not sure on how to add it to variadic arguments.

That's why I was asking if there will be a need either now or in the future -- is making this a parameter attribute the right design, or should this be a function attribute with positional arguments? (Not that I particularly love that design, but this isn't a user-facing attribute so if it's hard to use, I'm not that worried. I just want to make sure we don't need to add additional attributes in the future in this same space.)

[arsenm]

I can't see a scenario where we would need this with a variadic call. Ultimately these wrap specific physical instructions, not some kind of arbitrary call

[aaron.ballman]

FWIW, precommit CI has some relevant failures for a change:

******************** TEST 'Clang :: CodeGen/attr-maybeundef-template.cpp' FAILED ********************
Script:
--
: 'RUN: at line 1';   c:\ws\w4\llvm-project\premerge-checks\build\bin\clang.exe -cc1 -internal-isystem c:\ws\w4\llvm-project\premerge-checks\build\lib\clang\16.0.0\include -nostdsysteminc -no-opaque-pointers -emit-llvm C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp -o - | c:\ws\w4\llvm-project\premerge-checks\build\bin\filecheck.exe C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp
--
Exit Code: 1

Command Output (stdout):
--
$ ":" "RUN: at line 1"
$ "c:\ws\w4\llvm-project\premerge-checks\build\bin\clang.exe" "-cc1" "-internal-isystem" "c:\ws\w4\llvm-project\premerge-checks\build\lib\clang\16.0.0\include" "-nostdsysteminc" "-no-opaque-pointers" "-emit-llvm" "C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp" "-o" "-"
$ "c:\ws\w4\llvm-project\premerge-checks\build\bin\filecheck.exe" "C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp"
# command stderr:
C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp:3:11: error: CHECK: expected string not found in input
// CHECK: define weak_odr void @_Z5test4IfEvT_(float noundef [[TMP1:%.*]])
          ^
<stdin>:1:1: note: scanning from here
; ModuleID = 'C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp'
^
<stdin>:11:14: note: possible intended match here
define weak_odr dso_local void @"??$test4@M@@YAXM@Z"(float noundef %arg) #0 comdat {
             ^

Input file: <stdin>
Check file: C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           1: ; ModuleID = 'C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef-template.cpp' 
check:3'0     X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
           2: source_filename = "C:\\ws\\w4\\llvm-project\\premerge-checks\\clang\\test\\CodeGen\\attr-maybeundef-template.cpp" 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           3: target datalayout = "e-m:w-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128" 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           4: target triple = "x86_64-pc-windows-msvc" 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           5:  
check:3'0     ~
           6: $"??$test4@M@@YAXM@Z" = comdat any 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           7:  
check:3'0     ~
           8: $"??$test4@H@@YAXH@Z" = comdat any 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           9:  
check:3'0     ~
          10: ; Function Attrs: mustprogress noinline nounwind optnone 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          11: define weak_odr dso_local void @"??$test4@M@@YAXM@Z"(float noundef %arg) #0 comdat { 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
check:3'1                  ?                                                                        possible intended match
          12: entry: 
check:3'0     ~~~~~~~
          13:  %arg.addr = alloca float, align 4 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          14:  store float %arg, float* %arg.addr, align 4 
check:3'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          15:  ret void 
check:3'0     ~~~~~~~~~~
          16: } 
check:3'0     ~~
           .
           .
           .
>>>>>>

error: command failed with exit status: 1

--

********************
FAIL: Clang :: CodeGen/attr-maybeundef.c (4048 of 15750)
******************** TEST 'Clang :: CodeGen/attr-maybeundef.c' FAILED ********************
Script:
--
: 'RUN: at line 1';   c:\ws\w4\llvm-project\premerge-checks\build\bin\clang.exe -cc1 -internal-isystem c:\ws\w4\llvm-project\premerge-checks\build\lib\clang\16.0.0\include -nostdsysteminc -no-opaque-pointers -emit-llvm C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c -o - | c:\ws\w4\llvm-project\premerge-checks\build\bin\filecheck.exe C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c
--
Exit Code: 1

Command Output (stdout):
--
$ ":" "RUN: at line 1"
$ "c:\ws\w4\llvm-project\premerge-checks\build\bin\clang.exe" "-cc1" "-internal-isystem" "c:\ws\w4\llvm-project\premerge-checks\build\lib\clang\16.0.0\include" "-nostdsysteminc" "-no-opaque-pointers" "-emit-llvm" "C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c" "-o" "-"
# command stderr:
C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c:33:6: warning: 'maybe_undef' attribute only applies to parameters [-Wignored-attributes]

void __maybe_undef t2(int param1, int param2, float param3) {

     ^

C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c:3:38: note: expanded from macro '__maybe_undef'

#define __maybe_undef __attribute__((maybe_undef))

                                     ^

1 warning generated.


$ "c:\ws\w4\llvm-project\premerge-checks\build\bin\filecheck.exe" "C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c"
# command stderr:
C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c:7:16: error: CHECK-NEXT: is not on the line after the previous match
// CHECK-NEXT: [[TMP4:%.*]] = alloca i32, align 4
               ^
<stdin>:10:2: note: 'next' match was here
 %param2.addr = alloca i32, align 4
 ^
<stdin>:8:7: note: previous match ended here
entry:
      ^
<stdin>:9:1: note: non-matching line after previous match is here
 %param3.addr = alloca float, align 4
^

Input file: <stdin>
Check file: C:\ws\w4\llvm-project\premerge-checks\clang\test\CodeGen\attr-maybeundef.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
        .
        .
        .
        5:  
        6: ; Function Attrs: noinline nounwind optnone 
        7: define dso_local void @t1(i32 noundef %param1, i32 noundef %param2, float noundef %param3) #0 { 
        8: entry: 
        9:  %param3.addr = alloca float, align 4 
       10:  %param2.addr = alloca i32, align 4 
next:7      !~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~  error: match on wrong line
       11:  %param1.addr = alloca i32, align 4 
       12:  store float %param3, float* %param3.addr, align 4 
       13:  store i32 %param2, i32* %param2.addr, align 4 
       14:  store i32 %param1, i32* %param1.addr, align 4 
       15:  ret void 
        .
        .
        .
>>>>>>

error: command failed with exit status: 1

--

In D130224#3683831, @arsenm wrote:

I can't see a scenario where we would need this with a variadic call. Ultimately these wrap specific physical instructions, not some kind of arbitrary call

Excellent, that removes that concern. Thanks! Aside from the test failures and a minor coding style nit, I think this is close to ready to go.

clang/lib/CodeGen/CGCall.cpp

2063–2065

[skc7]

Fix tests failures.

[aaron.ballman]

clang/test/CodeGen/attr-maybeundef-template.cpp

4

This mangling works for Itanium targets, but doesn't match MSVC targets, which is why the test is failing on Windows.

clang/test/CodeGen/attr-maybeundef.c

49

It looks like dso_local is missing here on Windows and expected to be missing on Linux? Not certain what's up with that...

[skc7]

Fix windows tests failure

[skc7]

clang/test/CodeGen/attr-maybeundef.c

49

On Linux, dso_local isn't present.
declare void @VariadicFunction(i32 noundef, ...)

[aaron.ballman]

LGTM!

[nlopes]

Would it be possible to implement this by dropping the noundef attribute rather than emitting a freeze? Seems like a much better option to me.

[arsenm]

In D130224#3686756, @nlopes wrote:

Would it be possible to implement this by dropping the noundef attribute rather than emitting a freeze? Seems like a much better option to me.

See @nhaehnle's post here

[nlopes]

In D130224#3687351, @arsenm wrote:

In D130224#3686756, @nlopes wrote:

Would it be possible to implement this by dropping the noundef attribute rather than emitting a freeze? Seems like a much better option to me.

See @nhaehnle's post here

Thank you for the reference!
Seems like a big hammer, but I don't know the semantics of the GPU languages to comment further.

[mubarizafzal]

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

[amyk]

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

[aaron.ballman]

In D130224#3687860, @amyk wrote:

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

[skc7]

In D130224#3687907, @aaron.ballman wrote:

In D130224#3687860, @amyk wrote:

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Issue is with missing target triple in the tests. Submitted D130790 for review, which should fix the tests.

[amyk]

In D130224#3687907, @aaron.ballman wrote:

In D130224#3687860, @amyk wrote:

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Thanks @aaron.ballman. I've reverted the patch for the time being: https://reviews.llvm.org/rG4e1fe968c9de73507a1bf0c8aa57e06be457816e

[amyk]

In D130224#3688034, @skc7 wrote:

In D130224#3687907, @aaron.ballman wrote:

In D130224#3687860, @amyk wrote:

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230

Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Issue is with missing target triple in the tests. Submitted D130790 for review, which should fix the tests.

I realized I didn't happen to see this comment in time and had already reverted the patch. My apologies on this.

[skc7]

In D130224#3688067, @amyk wrote:

In D130224#3688034, @skc7 wrote:

In D130224#3687907, @aaron.ballman wrote:

In D130224#3687860, @amyk wrote:

In D130224#3687487, @mubarizafzal wrote:

Hi, the test cases that this patch introduces are failing on some ppc64le (Linux on Power) buildbots:
https://lab.llvm.org/buildbot/#/builders/57
https://lab.llvm.org/buildbot/#/builders/230
Would you mind taking a look please?

Thanks for pinging this patch. It appears it's affecting both little endian PPC:
https://lab.llvm.org/buildbot/#/builders/230/builds/1079
https://lab.llvm.org/buildbot/#/builders/121/builds/21978
https://lab.llvm.org/buildbot/#/builders/57/builds/20484
https://lab.llvm.org/buildbot/#/builders/36/builds/23702

And big endian PPC bots:
https://lab.llvm.org/buildbot/#/builders/231/builds/842
https://lab.llvm.org/buildbot/#/builders/93/builds/10270

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Issue is with missing target triple in the tests. Submitted D130790 for review, which should fix the tests.

I realized I didn't happen to see this comment in time and had already reverted the patch. My apologies on this.

Relanded the revert with fixes to tests. commit

clang/include/clang/Basic/Attr.td

 }
 
 def NoEscape : Attr {
   let Spellings = [Clang<"noescape">];
   let Subjects = SubjectList<[ParmVar]>;
   let Documentation = [NoEscapeDocs];
 }
 
+def MaybeUndef : InheritableAttr {
+  let Spellings = [Clang<"maybe_undef">];
+  let Subjects = SubjectList<[ParmVar]>;
+  let Documentation = [MaybeUndefDocs];
+  let SimpleHandler = 1;
+}
+
 def AssumeAligned : InheritableAttr {
   let Spellings = [GCC<"assume_aligned">];
   let Subjects = SubjectList<[ObjCMethod, Function]>;
   let Args = [ExprArgument<"Alignment">, ExprArgument<"Offset", 1>];
   let Documentation = [AssumeAlignedDocs];
 }
 
 def AllocAlign : InheritableAttr {

clang/include/clang/Basic/AttrDocs.td

.. code-block:: c
   void escapingFunc(__attribute__((noescape)) BlockTy block) {
     g0 = block; // Not OK.
     g1 = Block_copy(block); // Not OK either.
   }
 
   }];
 }
 
+def MaybeUndefDocs : Documentation {
+  let Category = DocCatVariable;
+  let Content = [{
+The ``maybe_undef`` attribute can be placed on a function parameter. It indicates
+that the parameter is allowed to use undef values. It informs the compiler
+to insert a freeze LLVM IR instruction on the function parameter.
+Please note that this is an attribute that is used as an internal
+implementation detail and not intended to be used by external users.
+
+In languages HIP, CUDA etc., some functions have multi-threaded semantics and
+it is enough for only one or some threads to provide defined arguments.
+Depending on semantics, undef arguments in some threads don't produce
+undefined results in the function call. Since, these functions accept undefined
+arguments, ``maybe_undef`` attribute can be placed.
+
+Sample usage:
+.. code-block:: c
+
+  void maybeundeffunc(int __attribute__((maybe_undef))param);
+  }];
+}
+
 def CarriesDependencyDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 The ``carries_dependency`` attribute specifies dependency propagation into and
 out of functions.
 
 When specified on a function or Objective-C method, the ``carries_dependency``
 attribute means that the return value carries a dependency out of the function,

clang/lib/CodeGen/CGCall.cpp

if (const MatrixType *Matrix = dyn_cast<MatrixType>(QTy))
     return DetermineNoUndef(Matrix->getElementType(), Types, DL, AI, false);
   if (const ArrayType *Array = dyn_cast<ArrayType>(QTy))
     return DetermineNoUndef(Array->getElementType(), Types, DL, AI, false);
 
   // TODO: Some structs may be `noundef`, in specific situations.
   return false;
 }
 
+/// Check if the argument of a function has maybe_undef attribute.
+static bool IsArgumentMaybeUndef(const Decl *TargetDecl,
+                                 unsigned NumRequiredArgs, unsigned ArgNo) {
+  const auto *FD = dyn_cast_or_null<FunctionDecl>(TargetDecl);
+  if (!FD)
+    return false;
+
+  // Assume variadic arguments do not have maybe_undef attribute.
+  if (ArgNo >= NumRequiredArgs)
+    return false;
+
+  // Check if argument has maybe_undef attribute.
+  if (ArgNo < FD->getNumParams()) {
+    const ParmVarDecl *Param = FD->getParamDecl(ArgNo);
+    if (Param && Param->hasAttr<MaybeUndefAttr>()) {
+      return true;
+    }
+  }
+
+  return false;
+}
+
 /// Construct the IR attribute list of a function or call.
 ///
 /// When adding an attribute, please consider where it should be handled:
 ///
 ///   - getDefaultFunctionAttributes is for attributes that are essentially
 ///     part of the global target configuration (but perhaps can be
 ///     overridden on a per-function basis).  Adding attributes there
 ///     will cause them to also be set in frontends that build on Clang's
for (CallArgList::const_iterator I = CallArgs.begin(), E = CallArgs.end();
     // Insert a padding argument to ensure proper alignment.
     if (IRFunctionArgs.hasPaddingArg(ArgNo))
       IRCallArgs[IRFunctionArgs.getPaddingArgNo(ArgNo)] =
           llvm::UndefValue::get(ArgInfo.getPaddingType());
 
     unsigned FirstIRArg, NumIRArgs;
     std::tie(FirstIRArg, NumIRArgs) = IRFunctionArgs.getIRArgs(ArgNo);
 
+    bool ArgHasMaybeUndefAttr =
+        IsArgumentMaybeUndef(TargetDecl, CallInfo.getNumRequiredArgs(), ArgNo);
+
     switch (ArgInfo.getKind()) {
     case ABIArgInfo::InAlloca: {
       assert(NumIRArgs == 0);
       assert(getTarget().getTriple().getArch() == llvm::Triple::x86);
       if (I->isAggregate()) {
         Address Addr = I->hasLValue()
                            ? I->getKnownLValue().getAddress(*this)
                            : I->getKnownRValue().getAggregateAddress();
for (CallArgList::const_iterator I = CallArgs.begin(), E = CallArgs.end();
 
     case ABIArgInfo::Indirect:
     case ABIArgInfo::IndirectAliased: {
       assert(NumIRArgs == 1);
       if (!I->isAggregate()) {
         // Make a temporary alloca to pass the argument.
         Address Addr = CreateMemTempWithoutCast(
             I->Ty, ArgInfo.getIndirectAlign(), "indirect-arg-temp");
-        IRCallArgs[FirstIRArg] = Addr.getPointer();
+
+        llvm::Value *Val = Addr.getPointer();
+        if (ArgHasMaybeUndefAttr)
+          Val = Builder.CreateFreeze(Addr.getPointer());
+        IRCallArgs[FirstIRArg] = Val;
 
         I->copyInto(*this, Addr);
       } else {
         // We want to avoid creating an unnecessary temporary+copy here;
         // however, we need one in three cases:
         // 1. If the argument is not byval, and we are required to copy the
         //    source.  (This case doesn't occur on any common architecture.)
         // 2. If the argument is byval, RV is not sufficiently aligned, and
case ABIArgInfo::IndirectAliased: {
             NeedCopy = true;
           }
         }
 
         if (NeedCopy) {
           // Create an aligned temporary, and copy to it.
           Address AI = CreateMemTempWithoutCast(
               I->Ty, ArgInfo.getIndirectAlign(), "byval-temp");
-          IRCallArgs[FirstIRArg] = AI.getPointer();
+          llvm::Value *Val = AI.getPointer();
+          if (ArgHasMaybeUndefAttr)
+            Val = Builder.CreateFreeze(AI.getPointer());
+          IRCallArgs[FirstIRArg] = Val;
 
           // Emit lifetime markers for the temporary alloca.
           llvm::TypeSize ByvalTempElementSize =
               CGM.getDataLayout().getTypeAllocSize(AI.getElementType());
           llvm::Value *LifetimeSize =
               EmitLifetimeStart(ByvalTempElementSize, AI.getPointer());
 
           // Add cleanup code to emit the end lifetime marker after the call.
           if (LifetimeSize) // In case we disabled lifetime markers.
             CallLifetimeEndAfterCall.emplace_back(AI, LifetimeSize);
 
           // Generate the copy.
           I->copyInto(*this, AI);
         } else {
           // Skip the extra memcpy call.
           auto *T = llvm::PointerType::getWithSamePointeeType(
               cast<llvm::PointerType>(V->getType()),
               CGM.getDataLayout().getAllocaAddrSpace());
-          IRCallArgs[FirstIRArg] = getTargetHooks().performAddrSpaceCast(
+
+          llvm::Value *Val = getTargetHooks().performAddrSpaceCast(
               *this, V, LangAS::Default, CGM.getASTAllocaAddressSpace(), T,
               true);
+          if (ArgHasMaybeUndefAttr)
+            Val = Builder.CreateFreeze(Val);
+          IRCallArgs[FirstIRArg] = Val;
         }
       }
       break;
     }
 
     case ABIArgInfo::Ignore:
       assert(NumIRArgs == 0);
       break;
case ABIArgInfo::Direct: {
           V = Builder.CreateZExt(V, ArgInfo.getCoerceToType());
 
         // If the argument doesn't match, perform a bitcast to coerce it.  This
         // can happen due to trivial type mismatches.
         if (FirstIRArg < IRFuncTy->getNumParams() &&
             V->getType() != IRFuncTy->getParamType(FirstIRArg))
           V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg));
 
+        if (ArgHasMaybeUndefAttr)
+          V = Builder.CreateFreeze(V);
         IRCallArgs[FirstIRArg] = V;
         break;
       }
 
       // FIXME: Avoid the conversion through memory if possible.
       Address Src = Address::invalid();
       if (!I->isAggregate()) {
         Src = CreateMemTemp(I->Ty, "coerce");
case ABIArgInfo::Direct: {
         } else {
           Src = Builder.CreateElementBitCast(Src, STy);
         }
 
         assert(NumIRArgs == STy->getNumElements());
         for (unsigned i = 0, e = STy->getNumElements(); i != e; ++i) {
           Address EltPtr = Builder.CreateStructGEP(Src, i);
           llvm::Value *LI = Builder.CreateLoad(EltPtr);
+          if (ArgHasMaybeUndefAttr)
+            LI = Builder.CreateFreeze(LI);
           IRCallArgs[FirstIRArg + i] = LI;
         }
       } else {
         // In the simple case, just pass the coerced loaded value.
         assert(NumIRArgs == 1);
         llvm::Value *Load =
             CreateCoercedLoad(Src, ArgInfo.getCoerceToType(), *this);
 
         if (CallInfo.isCmseNSCall()) {
           // For certain parameter types, clear padding bits, as they may reveal
           // sensitive information.
           // Small struct/union types are passed as integer arrays.
           auto *ATy = dyn_cast<llvm::ArrayType>(Load->getType());
           if (ATy != nullptr && isa<RecordType>(I->Ty.getCanonicalType()))
             Load = EmitCMSEClearRecord(Load, ATy, I->Ty);
         }
+
+        if (ArgHasMaybeUndefAttr)
+          Load = Builder.CreateFreeze(Load);
         IRCallArgs[FirstIRArg] = Load;
       }
 
       break;
     }
 
     case ABIArgInfo::CoerceAndExpand: {
       auto coercionType = ArgInfo.getCoerceAndExpandType();
case ABIArgInfo::CoerceAndExpand: {
       addr = Builder.CreateElementBitCast(addr, coercionType);
 
       unsigned IRArgPos = FirstIRArg;
       for (unsigned i = 0, e = coercionType->getNumElements(); i != e; ++i) {
         llvm::Type *eltType = coercionType->getElementType(i);
         if (ABIArgInfo::isPaddingForCoerceAndExpand(eltType)) continue;
         Address eltAddr = Builder.CreateStructGEP(addr, i);
         llvm::Value *elt = Builder.CreateLoad(eltAddr);
+        if (ArgHasMaybeUndefAttr)
+          elt = Builder.CreateFreeze(elt);
         IRCallArgs[IRArgPos++] = elt;
       }
       assert(IRArgPos == FirstIRArg + NumIRArgs);
 
       if (tempSize) {
         EmitLifetimeEnd(tempSize, AllocaAddr.getPointer());
       }
 

clang/lib/Sema/SemaDeclAttr.cpp

else
       handleNonNullAttr(S, D, AL);
     break;
   case ParsedAttr::AT_ReturnsNonNull:
     handleReturnsNonNullAttr(S, D, AL);
     break;
   case ParsedAttr::AT_NoEscape:
     handleNoEscapeAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_MaybeUndef:
+    handleSimpleAttribute<MaybeUndefAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_AssumeAligned:
     handleAssumeAlignedAttr(S, D, AL);
     break;
   case ParsedAttr::AT_AllocAlign:
     handleAllocAlignAttr(S, D, AL);
     break;
   case ParsedAttr::AT_Ownership:
     handleOwnershipAttr(S, D, AL);

clang/test/CodeGen/attr-maybeundef-template.cpp

+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
+
+// CHECK:      define weak_odr void @_Z5test4IfEvT_(float noundef [[TMP1:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca float, align 4
+// CHECK-NEXT:   store float [[TMP1:%.*]], float* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+// CHECK:      define weak_odr void @_Z5test4IiEvT_(i32 noundef [[TMP1:%.*]]) #0 comdat {
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+// CHECK:      define dso_local void @_Z4testv() #0 {
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP1:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca float, align 4
+// CHECK-NEXT:   [[TMP3:%.*]] = load i32, i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:   [[TMP4:%.*]] = freeze i32 [[TMP3:%.*]]
+// CHECK-NEXT:   call void @_Z5test4IiEvT_(i32 noundef [[TMP4:%.*]])
+// CHECK-NEXT:   [[TMP5:%.*]] = load float, float* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   [[TMP6:%.*]] = freeze float [[TMP5:%.*]]
+// CHECK-NEXT:   call void @_Z5test4IfEvT_(float noundef [[TMP6:%.*]])
+// CHECK-NEXT:   ret void
+
+template<class T>
+void test4(T __attribute__((maybe_undef)) arg) {
+  return;
+}
+
+template
+void test4<float>(float arg);
+
+template
+void test4<int>(int arg);
+
+void test() {
+    int Var1;
+    float Var2;
+    test4<int>(Var1);
+    test4<float>(Var2);
+}

clang/test/CodeGen/attr-maybeundef.c

+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
+
+#define __maybe_undef __attribute__((maybe_undef))
+
+// CHECK:      define dso_local void @t1(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], float noundef [[TMP3:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP4:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP5:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP6:%.*]] = alloca float, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT:   store i32 [[TMP2:%.*]], i32* [[TMP5:%.*]], align 4
+// CHECK-NEXT:   store float [[TMP3:%.*]], float* [[TMP6:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+// CHECK:      define dso_local void @t2(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], float noundef [[TMP3:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP4:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP5:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP6:%.*]] = alloca float, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT:   store i32 [[TMP2:%.*]], i32* [[TMP5:%.*]], align 4
+// CHECK-NEXT:   store float [[TMP3:%.*]], float* [[TMP6:%.*]], align 4
+// CHECK-NEXT:   [[TMP7:%.*]] = load i32, i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT:   [[TMP8:%.*]] = load i32, i32* [[TMP5:%.*]], align 4
+// CHECK-NEXT:   [[TMP9:%.*]] = load float, float* [[TMP6:%.*]], align 4
+// CHECK-NEXT:   [[TMP10:%.*]] = freeze i32 [[TMP8:%.*]]
+// CHECK-NEXT:   call void @t1(i32 noundef [[TMP7:%.*]], i32 noundef [[TMP10:%.*]], float noundef [[TMP9:%.*]])
+// CHECK-NEXT:   ret void
+
+void t1(int param1, int __maybe_undef param2, float param3) {}
+
+// expected-error {{'maybe_undef' attribute only applies to parameters [-Wignored-attributes]}}
+void __maybe_undef t2(int param1, int param2, float param3) {
+    t1(param1, param2, param3);
+}
+
+// CHECK:      define dso_local void @TestVariadicFunction(i32 noundef [[TMP0:%.*]], ...)
+// CHECK-NEXT: entry:
+// CHECK-NEXT:  [[TMP1:%.*]] = alloca i32, align 4
+// CHECK-NEXT:  [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT:  store i32 [[TMP0:%.*]], i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:  [[TMP3:%.*]] = load i32, i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:  [[TMP4:%.*]] = load i32, i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:  [[TMP5:%.*]] = load i32, i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:  [[TMP5:%.*]] = freeze i32 [[TMP2:%.*]]
+// CHECK-NEXT:  call void (i32, ...) @VariadicFunction(i32 noundef [[TMP6:%.*]], i32 noundef [[TMP4:%.*]], i32 noundef [[TMP5:%.*]])
+// CHECK-NEXT:  ret void
+
+// CHECK: declare void @VariadicFunction(i32 noundef, ...)
+
+void VariadicFunction(int __maybe_undef x, ...);
+void TestVariadicFunction(int x, ...) {
+  int Var;
+  return VariadicFunction(x, Var, Var);
+}
+
+// CHECK:      define dso_local void @other()
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP1:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP2:%.*]] = load i32, i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:   call void @func(i32 noundef [[TMP2:%.*]])
+// CHECK-NEXT:   [[TMP3:%.*]] = load i32, i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:   [[TMP4:%.*]] = freeze i32 [[TMP3:%.*]]
+// CHECK-NEXT:   call void @func1(i32 noundef [[TMP4:%.*]])
+// CHECK-NEXT:   ret void
+
+// CHECK:      define dso_local void @func(i32 noundef [[TMP1:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+// CHECK:      define dso_local void @func1(i32 noundef [[TMP1:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+void func(int param);
+void func1(int __maybe_undef param);
+
+void other() {
+  int Var;
+  func(Var);
+  func1(Var);
+}
+
+void func(__maybe_undef int param) {}
+void func1(int param) {}
+
+// CHECK:      define dso_local void @foo(i32 noundef [[TMP1:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   store i32 [[TMP1:%.*]], i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT:   ret void
+
+// CHECK:      define dso_local void @bar()
+// CHECK-NEXT: entry:
+// CHECK-NEXT:   [[TMP1:%.*]] = alloca i32, align 4
+// CHECK-NEXT:   [[TMP2:%.*]] = load i32, i32* [[TMP1:%.*]], align 4
+// CHECK-NEXT:   call void @foo(i32 noundef [[TMP2:%.*]])
+// CHECK-NEXT:   ret void
+
+void foo(__maybe_undef int param);
+void foo(int param) {}
+
+void bar() {
+  int Var;
+  foo(Var);
+}

clang/test/CodeGenHIP/maybe_undef-attr-verify.hip

+// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm  %s \
+// RUN:   -o - | FileCheck %s
+
+// CHECK: define dso_local amdgpu_kernel void @_Z13shufflekernelv()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP1:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[TMP2:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[TMP3:%.*]] = addrspacecast i32 addrspace(5)* [[TMP1:%.*]] to i32*
+// CHECK-NEXT:    [[TMP4:%.*]] = addrspacecast i32 addrspace(5)* [[TMP2:%.*]] to i32*
+// CHECK-NEXT:    [[TMP5:%.*]] = load i32, i32* [[TMP3:%.*]], align 4
+// CHECK-NEXT:    [[TMP6:%.*]] = freeze i32 [[TMP5:%.*]]
+// CHECK-NEXT:    %call = call noundef i32 @_Z11__shfl_synciii(i32 noundef [[TMP6:%.*]], i32 noundef 64, i32 noundef 0) #4
+// CHECK-NEXT:    store i32 %call, i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT:  ret void
+
+// CHECK: define linkonce_odr noundef i32 @_Z11__shfl_synciii(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], i32 noundef [[TMP3:%.*]])
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+#define __maybe_undef __attribute__((maybe_undef))
+#define WARP_SIZE 64
+
+static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE;
+
+__device__ static inline unsigned int __lane_id() {
+    return  __builtin_amdgcn_mbcnt_hi(
+        -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
+}
+
+__device__
+inline
+int __shfl_sync(int __maybe_undef var, int src_lane, int width = warpSize) {
+    int self = __lane_id();
+    int index = src_lane + (self & ~(width-1));
+    return __builtin_amdgcn_ds_bpermute(index<<2, var);
+}
+
+__global__ void
+shufflekernel()
+{
+    int t;
+    int res;
+    res = __shfl_sync(t, WARP_SIZE, 0);
+}

clang/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: InitPriority (SubjectMatchRule_variable)
 // CHECK-NEXT: InternalLinkage (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: LTOVisibilityPublic (SubjectMatchRule_record)
 // CHECK-NEXT: Leaf (SubjectMatchRule_function)
 // CHECK-NEXT: LoaderUninitialized (SubjectMatchRule_variable_is_global)
 // CHECK-NEXT: Lockable (SubjectMatchRule_record)
 // CHECK-NEXT: MIGServerRoutine (SubjectMatchRule_function, SubjectMatchRule_objc_method, SubjectMatchRule_block)
 // CHECK-NEXT: MSStruct (SubjectMatchRule_record)
+// CHECK-NEXT: MaybeUndef (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: MicroMips (SubjectMatchRule_function)
 // CHECK-NEXT: MinSize (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: MinVectorWidth (SubjectMatchRule_function)
 // CHECK-NEXT: Mips16 (SubjectMatchRule_function)
 // CHECK-NEXT: MipsLongCall (SubjectMatchRule_function)
 // CHECK-NEXT: MipsShortCall (SubjectMatchRule_function)
 // CHECK-NEXT: NSConsumed (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: NSConsumesSelf (SubjectMatchRule_objc_method)

clang/test/Sema/attr-maybeundef.c

+// RUN: %clang_cc1  -fsyntax-only -verify %s
+
+// Decl annotations.
+void f(int __attribute__((maybe_undef)) *a);
+void (*fp)(int __attribute__((maybe_undef)) handle);
+__attribute__((maybe_undef)) int i(); // expected-warning {{'maybe_undef' attribute only applies to parameters}}
+int __attribute__((maybe_undef)) a; // expected-warning {{'maybe_undef' attribute only applies to parameters}}
+int (* __attribute__((maybe_undef)) fpt)(char *); // expected-warning {{'maybe_undef' attribute only applies to parameters}}
+void h(int *a __attribute__((maybe_undef("RandomString")))); // expected-error {{'maybe_undef' attribute takes no arguments}}
+
+// Type annotations.
+int __attribute__((maybe_undef)) ta; // expected-warning {{'maybe_undef' attribute only applies to parameters}}
+
+// Typedefs.
+typedef int callback(char *) __attribute__((maybe_undef)); // expected-warning {{'maybe_undef' attribute only applies to parameters}}