Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[Clang] Undef attribute for global variables
Summary: [Clang] Attribute to allow defining undef global 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 provides an 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 is loosely based on the cuda __shared__ and opencl __local variable implementation which also produces undef global variables. Reviewers: kcc, rjmccall, rsmith, glider, vitalybuka, pcc, eugenis, vlad.tsyrklevich, jdoerfert, gregrodgers, jfb, aaron.ballman Reviewed By: rjmccall, aaron.ballman Subscribers: Anastasia, aaron.ballman, davidb, Quuxplusone, dexonsmith, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D74361
- Loading branch information
1 parent
1458bb9
commit c45eaea
Showing
13 changed files
with
237 additions
and
2 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,24 @@ | ||
// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s | ||
|
||
// CHECK: @tentative_attr_first = global i32 undef, align 4 | ||
int tentative_attr_first __attribute__((loader_uninitialized)); | ||
int tentative_attr_first; | ||
|
||
// CHECK: @tentative_attr_second = global i32 undef, align 4 | ||
int tentative_attr_second; | ||
int tentative_attr_second __attribute__((loader_uninitialized)); | ||
|
||
// CHECK: @array = global [16 x float] undef, align 16 | ||
float array[16] __attribute__((loader_uninitialized)); | ||
|
||
typedef struct | ||
{ | ||
int x; | ||
float y; | ||
} s; | ||
|
||
// CHECK: @i = global %struct.s undef, align 4 | ||
s i __attribute__((loader_uninitialized)); | ||
|
||
// CHECK: @private_extern_ok = hidden global i32 undef, align 4 | ||
__private_extern__ int private_extern_ok __attribute__((loader_uninitialized)); |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,29 @@ | ||
// RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s | ||
|
||
// CHECK: @defn = global i32 undef | ||
int defn [[clang::loader_uninitialized]]; | ||
|
||
// CHECK: @_ZL11defn_static = internal global i32 undef | ||
static int defn_static [[clang::loader_uninitialized]] __attribute__((used)); | ||
|
||
// CHECK: @_ZZ4funcvE4data = internal global i32 undef | ||
int* func(void) | ||
{ | ||
static int data [[clang::loader_uninitialized]]; | ||
return &data; | ||
} | ||
|
||
class trivial | ||
{ | ||
float x; | ||
}; | ||
|
||
// CHECK: @ut = global %class.trivial undef | ||
trivial ut [[clang::loader_uninitialized]]; | ||
|
||
// CHECK: @arr = global [32 x double] undef, align 16 | ||
double arr[32] __attribute__((loader_uninitialized)); | ||
|
||
// Defining as arr2[] [[clang..]] raises the error: attribute cannot be applied to types | ||
// CHECK: @arr2 = global [4 x double] undef, align 16 | ||
double arr2 [[clang::loader_uninitialized]] [4]; |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,37 @@ | ||
// RUN: %clang_cc1 %s -verify -fsyntax-only | ||
// See also attr-loader-uninitialized.cpp | ||
|
||
int good __attribute__((loader_uninitialized)); | ||
static int local_ok __attribute__((loader_uninitialized)); | ||
int hidden_ok __attribute__((visibility("hidden"))) __attribute__((loader_uninitialized)); | ||
|
||
const int can_still_be_const __attribute__((loader_uninitialized)); | ||
|
||
extern int external_rejected __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{variable 'external_rejected' cannot be declared both 'extern' and with the 'loader_uninitialized' attribute}} | ||
|
||
int noargs __attribute__((loader_uninitialized(0))); | ||
// expected-error@-1 {{'loader_uninitialized' attribute takes no arguments}} | ||
|
||
int init_rejected __attribute__((loader_uninitialized)) = 42; | ||
// expected-error@-1 {{variable with 'loader_uninitialized' attribute cannot have an initializer}} | ||
|
||
int declaration_then_uninit_ok; | ||
int declaration_then_uninit_ok __attribute__((loader_uninitialized)); | ||
|
||
int definition_then_uninit_rejected = 0; | ||
int definition_then_uninit_rejected __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{redeclaration cannot add 'loader_uninitialized' attribute}} | ||
// expected-note@-3 {{previous definition is here}} | ||
|
||
int tentative_repeated_ok __attribute__((loader_uninitialized)); | ||
int tentative_repeated_ok __attribute__((loader_uninitialized)); | ||
|
||
__private_extern__ int private_extern_can_be_initialised = 10; | ||
__private_extern__ int therefore_uninit_private_extern_ok __attribute__((loader_uninitialized)); | ||
|
||
__private_extern__ int initialized_private_extern_rejected __attribute__((loader_uninitialized)) = 5; | ||
// expected-error@-1 {{variable with 'loader_uninitialized' attribute cannot have an initializer}} | ||
|
||
extern __attribute__((visibility("hidden"))) int extern_hidden __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{variable 'extern_hidden' cannot be declared both 'extern' and with the 'loader_uninitialized' attribute}} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,60 @@ | ||
// RUN: %clang_cc1 %s -verify -fsyntax-only | ||
|
||
int good __attribute__((loader_uninitialized)); | ||
static int local_ok __attribute__((loader_uninitialized)); | ||
int hidden_ok __attribute__((visibility("hidden"))) __attribute__((loader_uninitialized)); | ||
|
||
const int still_cant_be_const __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{default initialization of an object of const type}} | ||
extern int external_rejected __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{variable 'external_rejected' cannot be declared both 'extern' and with the 'loader_uninitialized' attribute}} | ||
|
||
int noargs __attribute__((loader_uninitialized(0))); | ||
// expected-error@-1 {{'loader_uninitialized' attribute takes no arguments}} | ||
|
||
int init_rejected __attribute__((loader_uninitialized)) = 42; | ||
// expected-error@-1 {{variable with 'loader_uninitialized' attribute cannot have an initializer}} | ||
|
||
void func() __attribute__((loader_uninitialized)) | ||
// expected-warning@-1 {{'loader_uninitialized' attribute only applies to global variables}} | ||
{ | ||
int local __attribute__((loader_uninitialized)); | ||
// expected-warning@-1 {{'loader_uninitialized' attribute only applies to global variables}} | ||
|
||
static int sl __attribute__((loader_uninitialized)); | ||
} | ||
|
||
struct s { | ||
__attribute__((loader_uninitialized)) int field; | ||
// expected-warning@-1 {{'loader_uninitialized' attribute only applies to global variables}} | ||
|
||
static __attribute__((loader_uninitialized)) int sfield; | ||
|
||
} __attribute__((loader_uninitialized)); | ||
// expected-warning@-1 {{'loader_uninitialized' attribute only applies to global variables}} | ||
|
||
int redef_attr_first __attribute__((loader_uninitialized)); | ||
int redef_attr_first; | ||
// expected-error@-1 {{redefinition of 'redef_attr_first'}} | ||
// expected-note@-3 {{previous definition is here}} | ||
|
||
int redef_attr_second; | ||
int redef_attr_second __attribute__((loader_uninitialized)); | ||
// expected-warning@-1 {{attribute declaration must precede definition}} | ||
// expected-note@-3 {{previous definition is here}} | ||
// expected-error@-3 {{redefinition of 'redef_attr_second'}} | ||
// expected-note@-5 {{previous definition is here}} | ||
|
||
struct trivial {}; | ||
|
||
trivial default_ok __attribute__((loader_uninitialized)); | ||
trivial value_rejected __attribute__((loader_uninitialized)) {}; | ||
// expected-error@-1 {{variable with 'loader_uninitialized' attribute cannot have an initializer}} | ||
|
||
struct nontrivial | ||
{ | ||
nontrivial() {} | ||
}; | ||
|
||
nontrivial needs_trivial_ctor __attribute__((loader_uninitialized)); | ||
// expected-error@-1 {{variable with 'loader_uninitialized' attribute must have a trivial default constructor}} |