[PATCH] D15305: [CUDA] Do not allow dynamic initialization of global device side variables.

[Artem Belevich via cfe-commits]

Richard,

On Fri, Jan 15, 2016 at 5:32 PM, Richard Smith <richard at metafoo.co.uk>
wrote:

> On Fri, Jan 15, 2016 at 5:29 PM, Richard Smith <richard at metafoo.co.uk>
> wrote:
> > On Fri, Jan 15, 2016 at 4:22 PM, Artem Belevich <tra at google.com> wrote:
> >> tra added inline comments.
> >>
> >> ================
> >> Comment at: lib/CodeGen/CodeGenModule.cpp:2334
> >> @@ -2339,1 +2333,3 @@
> >> +      D->hasAttr<CUDASharedAttr>())
> >>      Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
> >> +  else if (!InitExpr) {
> >> ----------------
> >> rsmith wrote:
> >>> As this is a global variable, it should presumably still be statically
> zero-initialized.
> >> There is no way to initialize __shared__ variables. They are rough
> equivalent of local variables, only in this case CUDA allocates them per
> kernel invocation from a shared buffer with no guarantees regarding its
> contents.
> >>
> >> They used to be zero-initialized by compiler, but that was
> intentionally changed to undef in r245786 / http://reviews.llvm.org/D12241
> >
> > That doesn't seem right. C++ guarantees zero-initialization for all
> > globals, prior to performing any other initialization.
>
> It looks like the problem being fixed by D12241 was probably caused by
> the __shared__ variables having the wrong linkage.
>

I'll take a look at this separately as it's unrelated to this patch.

I believe current patch addresses your other comments.

--Artem




-- 
--Artem Belevich
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160201/d1e1f2d6/attachment.html>