JonChesterfield created this revision.
JonChesterfield added reviewers: kcc, rjmccall, rsmith, glider, vitalybuka,
pcc, eugenis, vlad.tsyrklevich, jdoerfert, gregrodgers.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
[Clang] Uninitialize attribute on global variables
Extends D54604 <https://reviews.llvm.org/D54604> to permit [[uninitialized]] on
global and static variables
Initializing global variables is very cheap on hosted implementations. The
C semantics of zero initializing globals work very well there. It is not
necessarily cheap on freestanding implementations. Where there is no loader
available, code must be emitted near the start point to write the appropriate
values into memory.
At present, external variables can be declared in C++ and definitions provided
in assembly (or IR) to achive this effect. This patch removes a restriction on
the existing attribute in order to remove this reason for writing assembly for
performance sensitive freestanding implementations.
A close analogue in tree is LDS memory for amdgcn, where the kernel is
responsible for initializing the memory after it starts executing on the gpu.
Uninitalized variables in LDS are observably cheaper than zero initialized.
Patch follows the cuda __shared__ variable implementation which also produces
undef global variables, and reuses the [[uninitialized]] attribute from auto
variable initialisation. I think the existing docs are still appropriate.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D74361
Files:
clang/include/clang/Basic/Attr.td
clang/lib/CodeGen/CGDecl.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/CodeGenCXX/attribute_uninitialized.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/test/Sema/attr-uninitialized.c
Index: clang/test/Sema/attr-uninitialized.c
===================================================================
--- clang/test/Sema/attr-uninitialized.c
+++ clang/test/Sema/attr-uninitialized.c
@@ -6,16 +6,11 @@
void bad() {
int im_bad __attribute((uninitialized("zero"))); // expected-error {{'uninitialized' attribute takes no arguments}}
- static int im_baaad __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to local variables}}
}
-extern int come_on __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to local variables}}
-int you_know __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to local variables}}
-static int and_the_whole_world_has_to __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to local variables}}
-
-void answer_right_now() __attribute((uninitialized)) {} // expected-warning {{'uninitialized' attribute only applies to local variables}}
-void just_to_tell_you_once_again(__attribute((uninitialized)) int whos_bad) {} // expected-warning {{'uninitialized' attribute only applies to local variables}}
+void answer_right_now() __attribute((uninitialized)) {} // expected-warning {{'uninitialized' attribute only applies to variables}}
+void just_to_tell_you_once_again(__attribute((uninitialized)) int whos_bad) {}
struct TheWordIsOut {
- __attribute((uninitialized)) int youre_doin_wrong; // expected-warning {{'uninitialized' attribute only applies to local variables}}
-} __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to local variables}}
+ __attribute((uninitialized)) int youre_doin_wrong; // expected-warning {{'uninitialized' attribute only applies to variables}}
+} __attribute((uninitialized)); // expected-warning {{'uninitialized' attribute only applies to variables}}
Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test
===================================================================
--- clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -148,6 +148,7 @@
// CHECK-NEXT: Target (SubjectMatchRule_function)
// CHECK-NEXT: TestTypestate (SubjectMatchRule_function_is_member)
// CHECK-NEXT: TrivialABI (SubjectMatchRule_record)
+// CHECK-NEXT: Uninitialized (SubjectMatchRule_variable)
// CHECK-NEXT: UseHandle (SubjectMatchRule_variable_is_parameter)
// CHECK-NEXT: VecReturn (SubjectMatchRule_record)
// CHECK-NEXT: VecTypeHint (SubjectMatchRule_function)
Index: clang/test/CodeGenCXX/attribute_uninitialized.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenCXX/attribute_uninitialized.cpp
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -std=c++11 -emit-llvm -o - %s | FileCheck %s
+
+// CHECK: @_ZZ4funcvE4data = internal global i32 undef
+int* func(void)
+{
+ static int data [[clang::uninitialized]];
+ return &data;
+}
+
+// No code emitted
+extern int extern_unhelpful_but_harmless [[clang::uninitialized]];
+
+// CHECK: @tentative = global i32 undef
+int tentative [[clang::uninitialized]];
+
+// CHECK: @_ZL16tentative_static = internal global i32 undef
+static int tentative_static [[clang::uninitialized]] __attribute__((used));
+
+// CHECK: @nominally_zero_init = global i32 undef
+int nominally_zero_init [[clang::uninitialized]] = 0;
+
+// CHECK: @nominally_value_init = global i32 undef
+int nominally_value_init [[clang::uninitialized]] = 4;
+
+class trivial
+{
+ float x;
+};
+
+// CHECK: @ut = global %class.trivial undef
+trivial ut [[clang::uninitialized]]; // claims zero init
+
+struct nontrivial
+{
+ nontrivial() : x(3.14) {}
+ double x;
+};
+
+// CHECK: @unt = global %struct.nontrivial undef
+nontrivial unt [[clang::uninitialized]];
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -6501,8 +6501,6 @@
}
static void handleUninitializedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
- assert(cast<VarDecl>(D)->getStorageDuration() == SD_Automatic &&
- "uninitialized is only valid on automatic duration variables");
D->addAttr(::new (S.Context) UninitializedAttr(S.Context, AL));
}
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3943,9 +3943,12 @@
// left undefined.
bool IsHIPPinnedShadowVar =
getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
+ bool IsUninitializedVar = D->hasAttr<UninitializedAttr>();
if (getLangOpts().CUDA &&
(IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
+ else if (IsUninitializedVar)
+ Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
// implicitly initialized with { 0 }.
Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ clang/lib/CodeGen/CGDecl.cpp
@@ -249,7 +249,8 @@
// variables cannot have an initializer.
llvm::Constant *Init = nullptr;
if (Ty.getAddressSpace() == LangAS::opencl_local ||
- D.hasAttr<CUDASharedAttr>())
+ D.hasAttr<CUDASharedAttr>() ||
+ D.hasAttr<UninitializedAttr>())
Init = llvm::UndefValue::get(LTy);
else
Init = EmitNullConstant(Ty);
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -3473,7 +3473,7 @@
def Uninitialized : InheritableAttr {
let Spellings = [Clang<"uninitialized", 0>];
- let Subjects = SubjectList<[LocalVar]>;
+ let Subjects = SubjectList<[Var]>;
let Documentation = [UninitializedDocs];
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
