c45eaeabb7
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
30 lines
832 B
C++
30 lines
832 B
C++
// 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];
|