maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

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

This reverts commit a35c64ce23.

Reverting this commit as it causes various failures on LE and BE PPC bots.
This commit is contained in:
Amy Kwan 2022-07-29 13:16:37 -05:00
parent cea1b790f6
commit 4e1fe968c9
9 changed files with 3 additions and 291 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
clang/include/clang/Basic/Attr.td
View File

@ -2023,13 +2023,6 @@ 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]>;

22
clang/include/clang/Basic/AttrDocs.td
View File

@ -257,28 +257,6 @@ 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 = [{

50
clang/lib/CodeGen/CGCall.cpp
View File

@ -2046,27 +2046,6 @@ 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:
@ -4842,9 +4821,6 @@ 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);
@ -4903,11 +4879,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
// Make a temporary alloca to pass the argument.
Address Addr = CreateMemTempWithoutCast(
I->Ty, ArgInfo.getIndirectAlign(), "indirect-arg-temp");
llvm::Value *Val = Addr.getPointer();
if (ArgHasMaybeUndefAttr)
Val = Builder.CreateFreeze(Addr.getPointer());
IRCallArgs[FirstIRArg] = Val;
IRCallArgs[FirstIRArg] = Addr.getPointer();
I->copyInto(*this, Addr);
} else {
@ -4965,10 +4937,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
// Create an aligned temporary, and copy to it.
Address AI = CreateMemTempWithoutCast(
I->Ty, ArgInfo.getIndirectAlign(), "byval-temp");
llvm::Value *Val = AI.getPointer();
if (ArgHasMaybeUndefAttr)
Val = Builder.CreateFreeze(AI.getPointer());
IRCallArgs[FirstIRArg] = Val;
IRCallArgs[FirstIRArg] = AI.getPointer();
// Emit lifetime markers for the temporary alloca.
llvm::TypeSize ByvalTempElementSize =
@ -4987,13 +4956,9 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
auto *T = llvm::PointerType::getWithSamePointeeType(
cast<llvm::PointerType>(V->getType()),
CGM.getDataLayout().getAllocaAddrSpace());
llvm::Value *Val = getTargetHooks().performAddrSpaceCast(
IRCallArgs[FirstIRArg] = getTargetHooks().performAddrSpaceCast(
*this, V, LangAS::Default, CGM.getASTAllocaAddressSpace(), T,
true);
if (ArgHasMaybeUndefAttr)
Val = Builder.CreateFreeze(Val);
IRCallArgs[FirstIRArg] = Val;
}
}
break;
@ -5047,8 +5012,6 @@ 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;
}
@ -5093,8 +5056,6 @@ 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 {
@ -5111,9 +5072,6 @@ 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;
}
@ -5159,8 +5117,6 @@ 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);

3
clang/lib/Sema/SemaDeclAttr.cpp
View File

@ -8634,9 +8634,6 @@ 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;

43
clang/test/CodeGen/attr-maybeundef-template.cpp
View File

@ -1,43 +0,0 @@
// 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);
}

109
clang/test/CodeGen/attr-maybeundef.c
View File

@ -1,109 +0,0 @@
// 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);
}

44
clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
View File

@ -1,44 +0,0 @@
// 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);
}

1
clang/test/Misc/pragma-attribute-supported-attributes-list.test
View File

@ -83,7 +83,6 @@
// 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)

15
clang/test/Sema/attr-maybeundef.c
View File

@ -1,15 +0,0 @@
// 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}}