[clang] a35c64c - [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values
via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 28 19:38:55 PDT 2022
Author: skc7
Date: 2022-07-29T02:27:26Z
New Revision: a35c64ce23b7c7e4972c89b224b9363639dddea2
URL: https://github.com/llvm/llvm-project/commit/a35c64ce23b7c7e4972c89b224b9363639dddea2
DIFF: https://github.com/llvm/llvm-project/commit/a35c64ce23b7c7e4972c89b224b9363639dddea2.diff
LOG: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values
Add the ability to put __attribute__((maybe_undef)) on function arguments.
Clang codegen introduces a freeze instruction on the argument.
Differential Revision: https://reviews.llvm.org/D130224
Added:
clang/test/CodeGen/attr-maybeundef-template.cpp
clang/test/CodeGen/attr-maybeundef.c
clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
clang/test/Sema/attr-maybeundef.c
Modified:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/lib/CodeGen/CGCall.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 0460371d26c94..a94829698ad91 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2023,6 +2023,13 @@ def NoEscape : Attr {
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]>;
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 5c84e2fc5b77d..f61a5a8d5b523 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -257,6 +257,28 @@ applies to copies of the block. For example:
}];
}
+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 = [{
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 7853695f1f0cb..ee37e762dc759 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -2046,6 +2046,27 @@ static bool DetermineNoUndef(QualType QTy, CodeGenTypes &Types,
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:
@@ -4821,6 +4842,9 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
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);
@@ -4879,7 +4903,11 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
// 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 {
@@ -4937,7 +4965,10 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
// 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 =
@@ -4956,9 +4987,13 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
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;
@@ -5012,6 +5047,8 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
V->getType() != IRFuncTy->getParamType(FirstIRArg))
V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg));
+ if (ArgHasMaybeUndefAttr)
+ V = Builder.CreateFreeze(V);
IRCallArgs[FirstIRArg] = V;
break;
}
@@ -5056,6 +5093,8 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
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 {
@@ -5072,6 +5111,9 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
if (ATy != nullptr && isa<RecordType>(I->Ty.getCanonicalType()))
Load = EmitCMSEClearRecord(Load, ATy, I->Ty);
}
+
+ if (ArgHasMaybeUndefAttr)
+ Load = Builder.CreateFreeze(Load);
IRCallArgs[FirstIRArg] = Load;
}
@@ -5117,6 +5159,8 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
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);
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 838fd48357fb7..695fedc889fda 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -8634,6 +8634,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
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;
diff --git a/clang/test/CodeGen/attr-maybeundef-template.cpp b/clang/test/CodeGen/attr-maybeundef-template.cpp
new file mode 100644
index 0000000000000..33a999143546d
--- /dev/null
+++ b/clang/test/CodeGen/attr-maybeundef-template.cpp
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
+
+// CHECK-LABEL: define{{.*}} void @{{.*}}test4{{.*}}(float
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP2:%.*]] = alloca float, align 4
+// CHECK-NEXT: store float [[TMP1:%.*]], float* [[TMP2:%.*]], align 4
+// CHECK-NEXT: ret void
+
+// CHECK-LABEL: define{{.*}} void @{{.*}}test4{{.*}}(i32
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP2:%.*]] = alloca i32, align 4
+// CHECK-NEXT: store i32 [[TMP1:%.*]], i32* [[TMP2:%.*]], align 4
+// CHECK-NEXT: ret void
+
+// CHECK-LABEL: define{{.*}} void @{{.*}}test{{.*}}(
+// 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 @{{.*}}test4{{.*}}(i32 noundef [[TMP4:%.*]])
+// CHECK-NEXT: [[TMP5:%.*]] = load float, float* [[TMP2:%.*]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = freeze float [[TMP5:%.*]]
+// CHECK-NEXT: call void @{{.*}}test4{{.*}}(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);
+}
diff --git a/clang/test/CodeGen/attr-maybeundef.c b/clang/test/CodeGen/attr-maybeundef.c
new file mode 100644
index 0000000000000..b966ae998f77f
--- /dev/null
+++ b/clang/test/CodeGen/attr-maybeundef.c
@@ -0,0 +1,109 @@
+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
+
+#define __maybe_undef __attribute__((maybe_undef))
+
+// CHECK: define{{.*}} void @t1(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], i32 noundef [[TMP3:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP4:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[TMP5:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[TMP6:%.*]] = alloca i32, align 4
+// CHECK-NEXT: store i32 [[TMP1:%.*]], i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT: store i32 [[TMP2:%.*]], i32* [[TMP5:%.*]], align 4
+// CHECK-NEXT: store i32 [[TMP3:%.*]], i32* [[TMP6:%.*]], align 4
+// CHECK-NEXT: ret void
+
+// CHECK: define{{.*}} void @t2(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], i32 noundef [[TMP3:%.*]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP4:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[TMP5:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[TMP6:%.*]] = alloca i32, align 4
+// CHECK-NEXT: store i32 [[TMP1:%.*]], i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT: store i32 [[TMP2:%.*]], i32* [[TMP5:%.*]], align 4
+// CHECK-NEXT: store i32 [[TMP3:%.*]], i32* [[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 i32, i32* [[TMP6:%.*]], align 4
+// CHECK-NEXT: [[TMP10:%.*]] = freeze i32 [[TMP8:%.*]]
+// CHECK-NEXT: call void @t1(i32 noundef [[TMP7:%.*]], i32 noundef [[TMP10:%.*]], i32 noundef [[TMP9:%.*]])
+// CHECK-NEXT: ret void
+
+void t1(int param1, int __maybe_undef param2, int param3) {}
+
+void t2(int param1, int param2, int param3) {
+ t1(param1, param2, param3);
+}
+
+// CHECK: define{{.*}} 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{{.*}} 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{{.*}} 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{{.*}} 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{{.*}} 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{{.*}} 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);
+}
diff --git a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
new file mode 100644
index 0000000000000..afa461f909529
--- /dev/null
+++ b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
@@ -0,0 +1,44 @@
+// 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);
+}
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 64e2bf619004e..d6e1538bd92a1 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -83,6 +83,7 @@
// 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)
diff --git a/clang/test/Sema/attr-maybeundef.c b/clang/test/Sema/attr-maybeundef.c
new file mode 100644
index 0000000000000..09bb287bfc4b4
--- /dev/null
+++ b/clang/test/Sema/attr-maybeundef.c
@@ -0,0 +1,15 @@
+// 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}}
More information about the cfe-commits
mailing list