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

[David Blaikie]

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

> Added better explanation of the test.
>
>
> http://reviews.llvm.org/D9327
>
> Files:
>   lib/Sema/SemaDecl.cpp
>   test/SemaCUDA/qualifiers.cu
>
> Index: lib/Sema/SemaDecl.cpp
> ===================================================================
> --- lib/Sema/SemaDecl.cpp
> +++ lib/Sema/SemaDecl.cpp
> @@ -5769,12 +5769,16 @@
>             diag::err_thread_non_global)
>          << DeclSpec::getSpecifierName(TSCS);
>      else if (!Context.getTargetInfo().isTLSSupported()) {
> -      if (getLangOpts().CUDA)
> +      if (getLangOpts().CUDA) {
>          // Postpone error emission until we've collected attributes
> required to
>          // figure out whether it's a host or device variable and whether
> the
>          // error should be ignored.
>          EmitTLSUnsupportedError = true;
> -      else
> +        // We still need to mark the variable as TLS so it shows up in
> AST with
> +        // proper storage class for other tools to use even if we're not
> going
> +        // to emit any code for it.
> +        NewVD->setTSCSpec(TSCS);
> +      } else
>          Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
>               diag::err_thread_unsupported);
>      } else
> Index: test/SemaCUDA/qualifiers.cu
> ===================================================================
> --- 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)))


> +// 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
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20150428/288c73b8/attachment.html>