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