On Fri, Jan 15, 2016 at 5:29 PM, Richard Smith <rich...@metafoo.co.uk> wrote:
> On Fri, Jan 15, 2016 at 4:22 PM, Artem Belevich <t...@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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
