yaxunl created this revision. yaxunl added reviewers: tra, rjmccall. CUDA 8.0 E.3.9.4 says: Within the body of a __device__ or __global__ function, only __shared__ variables or variables without any device memory qualifiers may be declared with static storage class.
It is unclear how a function-scope non-const static variable without device memory qualifier is implemented, therefore only static const variable without device memory qualifier is allowed, which can be emitted as a global variable in constant address space. Currently clang only allows function-scope static variable with `__shared__` qualifier. This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space. https://reviews.llvm.org/D49931 Files: lib/CodeGen/CodeGenModule.cpp lib/Sema/SemaDecl.cpp test/CodeGenCUDA/device-var-init.cu Index: test/CodeGenCUDA/device-var-init.cu =================================================================== --- test/CodeGenCUDA/device-var-init.cu +++ test/CodeGenCUDA/device-var-init.cu @@ -112,6 +112,9 @@ // CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef // CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef +// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5] +// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123 + // We should not emit global initializers for device-side variables. // CHECK-NOT: @__cxx_global_var_init @@ -234,6 +237,9 @@ static __shared__ ETC s_etc; // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) + static const int const_array[] = {1, 2, 3, 4, 5}; + static const int const_int = 123; + // anchor point separating constructors and destructors df(); // CHECK: call void @_Z2dfv() Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -11914,10 +11914,15 @@ NewAttr->setInherited(true); VD->addAttr(NewAttr); } - // CUDA E.2.9.4: Within the body of a __device__ or __global__ - // function, only __shared__ variables may be declared with - // static storage class. - if (getLangOpts().CUDA && !VD->hasAttr<CUDASharedAttr>() && + // CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__ + // function, only __shared__ variables or variables without any device + // memory qualifiers may be declared with static storage class. + // Note: It is unclear how a function-scope non-const static variable + // without device memory qualifier is implemented, therefore only static + // const variable without device memory qualifier is allowed. + if (getLangOpts().CUDA && + !(VD->hasAttr<CUDASharedAttr>() || + VD->getType().isConstQualified()) && CUDADiagIfDeviceCode(VD->getLocation(), diag::err_device_static_local_var) << CurrentCUDATarget()) Index: lib/CodeGen/CodeGenModule.cpp =================================================================== --- lib/CodeGen/CodeGenModule.cpp +++ lib/CodeGen/CodeGenModule.cpp @@ -3176,6 +3176,10 @@ return LangAS::cuda_constant; else if (D && D->hasAttr<CUDASharedAttr>()) return LangAS::cuda_shared; + else if (D && D->hasAttr<CUDADeviceAttr>()) + return LangAS::cuda_device; + else if (D && D->getType().isConstQualified()) + return LangAS::cuda_constant; else return LangAS::cuda_device; }
Index: test/CodeGenCUDA/device-var-init.cu =================================================================== --- test/CodeGenCUDA/device-var-init.cu +++ test/CodeGenCUDA/device-var-init.cu @@ -112,6 +112,9 @@ // CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef // CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef +// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5] +// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123 + // We should not emit global initializers for device-side variables. // CHECK-NOT: @__cxx_global_var_init @@ -234,6 +237,9 @@ static __shared__ ETC s_etc; // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) + static const int const_array[] = {1, 2, 3, 4, 5}; + static const int const_int = 123; + // anchor point separating constructors and destructors df(); // CHECK: call void @_Z2dfv() Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -11914,10 +11914,15 @@ NewAttr->setInherited(true); VD->addAttr(NewAttr); } - // CUDA E.2.9.4: Within the body of a __device__ or __global__ - // function, only __shared__ variables may be declared with - // static storage class. - if (getLangOpts().CUDA && !VD->hasAttr<CUDASharedAttr>() && + // CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__ + // function, only __shared__ variables or variables without any device + // memory qualifiers may be declared with static storage class. + // Note: It is unclear how a function-scope non-const static variable + // without device memory qualifier is implemented, therefore only static + // const variable without device memory qualifier is allowed. + if (getLangOpts().CUDA && + !(VD->hasAttr<CUDASharedAttr>() || + VD->getType().isConstQualified()) && CUDADiagIfDeviceCode(VD->getLocation(), diag::err_device_static_local_var) << CurrentCUDATarget()) Index: lib/CodeGen/CodeGenModule.cpp =================================================================== --- lib/CodeGen/CodeGenModule.cpp +++ lib/CodeGen/CodeGenModule.cpp @@ -3176,6 +3176,10 @@ return LangAS::cuda_constant; else if (D && D->hasAttr<CUDASharedAttr>()) return LangAS::cuda_shared; + else if (D && D->hasAttr<CUDADeviceAttr>()) + return LangAS::cuda_device; + else if (D && D->getType().isConstQualified()) + return LangAS::cuda_constant; else return LangAS::cuda_device; }
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits