On Tue, Apr 28, 2015 at 1:31 PM, Artem Belevich <[email protected]> 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 > [email protected] > http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits > >
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
