Author: tra Date: Thu Sep 10 12:26:58 2015 New Revision: 247307 URL: http://llvm.org/viewvc/llvm-project?rev=247307&view=rev Log: [CUDA] Allow trivial constructors as initializer for __shared__ variables.
Differential Revision: http://reviews.llvm.org/D12739 Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp cfe/trunk/test/CodeGenCUDA/address-spaces.cu Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=247307&r1=247306&r2=247307&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original) +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Sep 10 12:26:58 2015 @@ -2165,8 +2165,10 @@ void CodeGenModule::EmitGlobalVarDefinit 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) { Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=247307&r1=247306&r2=247307&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original) +++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Thu Sep 10 12:26:58 2015 @@ -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; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
