[PATCH] [cuda] Preserve TLS storage class of host variable during device-side compilation.

[David Blaikie]

On Tue, Apr 28, 2015 at 2:24 PM, Artem Belevich <tra at google.com> wrote:

>
>
> On Tue, Apr 28, 2015 at 1:47 PM, David Blaikie <dblaikie at gmail.com> wrote:
>
>> --- test/SemaCUDA/qualifiers.cu
>>> +++ test/SemaCUDA/qualifiers.cu
>>> @@ -1,21 +1,35 @@
>>>  // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only
>>> -verify %s
>>>  // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify
>>> -fcuda-is-device %s
>>> +//
>>> +// We run clang_cc1 with 'not' because source file contains
>>> +// intentional errors. CC1 failure is expected and must be ignored
>>> +// here. We're interested in what ends up in AST and that's what
>>> +// FileCheck verifies.
>>>
>>
>> Why are there errors? (the ast of code that is incorrect isn't really
>> guaranteed, so this might be an unstable/unreliable test - perhaps we could
>> split out the error cases from the successful cases (& I'd still like to
>> consider IRGen testing the successful cases, rather than checking the AST -
>> the former seems like a better contract for Clang to abide by (the AST is a
>> bit more of an implementation detail)))
>>
>>
>
> The source code is syntactically valid. TLS is not supported by NVPTX and
> that's what produces the errors.
> IRGen is not going to work here, because host variables will not generate
> any code during device-side compilation.
>
> I guess I could extract a subset of this test that does not produce any
> errors but which would still allow me to check correctness of AST in a
> subset of practically useful cases. It does not help with potential
> instability of AST, though.
>

So what's the observable behavior of this change? "we've lost information
which is useful during source analysis." - useful for what? Diagnostic
messages? Then could we -verify check one of those diagnostics?


>
> --Artem
>
>
>
>
>> +// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only
>>> -ast-dump %s \
>>> +// RUN:   | FileCheck %s --check-prefix=CHECK-ALL
>>> --check-prefix=CHECK-HOST
>>> +// RUN: not %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only
>>> -ast-dump -fcuda-is-device %s \
>>> +// RUN:   | FileCheck %s --check-prefix=CHECK-ALL
>>> --check-prefix=CHECK-DEVICE
>>>
>>>  #include "Inputs/cuda.h"
>>>
>>>  // Host (x86) supports TLS and device-side compilation should ignore
>>>  // host variables. No errors in either case.
>>>  int __thread host_tls_var;
>>> +// CHECK-ALL: host_tls_var 'int' tls
>>>
>>>  #if defined(__CUDA_ARCH__)
>>>  // NVPTX does not support TLS
>>>  __device__ int __thread device_tls_var; // expected-error
>>> {{thread-local storage is not supported for the current target}}
>>> +// CHECK-DEVICE: device_tls_var 'int' tls
>>>  __shared__ int __thread shared_tls_var; // expected-error
>>> {{thread-local storage is not supported for the current target}}
>>> +// CHECK-DEVICE: shared_tls_var 'int' tls
>>>  #else
>>>  // Device-side vars should not produce any errors during host-side
>>>  // compilation.
>>>  __device__ int __thread device_tls_var;
>>> +// CHECK-HOST: device_tls_var 'int' tls
>>>  __shared__ int __thread shared_tls_var;
>>> +// CHECK-HOST: shared_tls_var 'int' tls
>>>  #endif
>>>
>>>  __global__ void g1(int x) {}
>>>
>>> EMAIL PREFERENCES
>>>   http://reviews.llvm.org/settings/panel/emailpreferences/
>>>
>>> _______________________________________________
>>> cfe-commits mailing list
>>> cfe-commits at cs.uiuc.edu
>>> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
>>>
>>>
>>
>
>
> --
> --Artem Belevich
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20150428/bfe98f89/attachment.html>