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

[Artem Belevich]

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.

--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/4d8f88a1/attachment.html>