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}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
