Author: Amy Kwan Date: 2022-07-29T13:28:48-05:00 New Revision: 4e1fe968c9de73507a1bf0c8aa57e06be457816e
URL: https://github.com/llvm/llvm-project/commit/4e1fe968c9de73507a1bf0c8aa57e06be457816e DIFF: https://github.com/llvm/llvm-project/commit/4e1fe968c9de73507a1bf0c8aa57e06be457816e.diff LOG: Revert "[Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values" This reverts commit a35c64ce23b7c7e4972c89b224b9363639dddea2. Reverting this commit as it causes various failures on LE and BE PPC bots. Added: 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: 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 ################################################################################ diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a94829698ad91..0460371d26c94 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -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]>; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index f61a5a8d5b523..5c84e2fc5b77d 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -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 = [{ diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index ee37e762dc759..7853695f1f0cb 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -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); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 695fedc889fda..838fd48357fb7 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -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; diff --git a/clang/test/CodeGen/attr-maybeundef-template.cpp b/clang/test/CodeGen/attr-maybeundef-template.cpp deleted file mode 100644 index 33a999143546d..0000000000000 --- a/clang/test/CodeGen/attr-maybeundef-template.cpp +++ /dev/null @@ -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); -} diff --git a/clang/test/CodeGen/attr-maybeundef.c b/clang/test/CodeGen/attr-maybeundef.c deleted file mode 100644 index b966ae998f77f..0000000000000 --- a/clang/test/CodeGen/attr-maybeundef.c +++ /dev/null @@ -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); -} diff --git a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip deleted file mode 100644 index afa461f909529..0000000000000 --- a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip +++ /dev/null @@ -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); -} diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index d6e1538bd92a1..64e2bf619004e 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -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) diff --git a/clang/test/Sema/attr-maybeundef.c b/clang/test/Sema/attr-maybeundef.c deleted file mode 100644 index 09bb287bfc4b4..0000000000000 --- a/clang/test/Sema/attr-maybeundef.c +++ /dev/null @@ -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}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
