Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-02-20 10:07:02 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[CUDA] Allow trivial constructors as initializer for __shared__ variables.

Browse Source 
Differential Revision: http://reviews.llvm.org/D12739

llvm-svn: 247307
This commit is contained in:
Artem Belevich 2015-09-10 17:26:58 +00:00
parent 170fd9cb6e
commit da1851ca58
2 changed files with 18 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -2165,8 +2165,10 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) {
if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
&& D->hasAttr<CUDASharedAttr>()) {
if (InitExpr) {
Error(D->getLocation(),
"__shared__ variable cannot have an initialization.");
const auto *C = dyn_cast<CXXConstructExpr>(InitExpr);
if (C == nullptr || !C->getConstructor()->hasTrivialBody())
Error(D->getLocation(),
"__shared__ variable cannot have an initialization.");
}
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
} else if (!InitExpr) {

14
clang/test/CodeGenCUDA/address-spaces.cu
View File

@ -25,6 +25,8 @@ struct MyStruct {
// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
// CHECK: @b = addrspace(3) global float undef
// CHECK: @c = addrspace(3) global %struct.c undef
// CHECK @d = addrspace(3) global %struct.d undef
__device__ void foo() {
// CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
@ -117,3 +119,15 @@ __device__ int construct_shared_struct() {
return t.getData();
// CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
}
// Make sure we allow __shared__ structures with default or empty constructors.
struct c {
int i;
};
__shared__ struct c c;
struct d {
int i;
d() {}
};
__shared__ struct d d;