<div dir="ltr"><br><div class="gmail_extra"><br><div class="gmail_quote">On Tue, Apr 28, 2015 at 1:31 PM, Artem Belevich <span dir="ltr"><<a href="mailto:tra@google.com" target="_blank">tra@google.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Added better explanation of the test.<br>
<span class=""><br>
<br>
<a href="http://reviews.llvm.org/D9327" target="_blank">http://reviews.llvm.org/D9327</a><br>
<br>
Files:<br>
lib/Sema/SemaDecl.cpp<br>
test/SemaCUDA/<a href="http://qualifiers.cu" target="_blank">qualifiers.cu</a><br>
<br>
</span><span class="">Index: lib/Sema/SemaDecl.cpp<br>
===================================================================<br>
--- lib/Sema/SemaDecl.cpp<br>
+++ lib/Sema/SemaDecl.cpp<br>
@@ -5769,12 +5769,16 @@<br>
diag::err_thread_non_global)<br>
<< DeclSpec::getSpecifierName(TSCS);<br>
else if (!Context.getTargetInfo().isTLSSupported()) {<br>
- if (getLangOpts().CUDA)<br>
+ if (getLangOpts().CUDA) {<br>
// Postpone error emission until we've collected attributes required to<br>
// figure out whether it's a host or device variable and whether the<br>
// error should be ignored.<br>
EmitTLSUnsupportedError = true;<br>
- else<br>
+ // We still need to mark the variable as TLS so it shows up in AST with<br>
+ // proper storage class for other tools to use even if we're not going<br>
+ // to emit any code for it.<br>
+ NewVD->setTSCSpec(TSCS);<br>
+ } else<br>
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),<br>
diag::err_thread_unsupported);<br>
} else<br>
</span><span class="">Index: test/SemaCUDA/<a href="http://qualifiers.cu" target="_blank">qualifiers.cu</a><br>
===================================================================<br>
--- test/SemaCUDA/<a href="http://qualifiers.cu" target="_blank">qualifiers.cu</a><br>
+++ test/SemaCUDA/<a href="http://qualifiers.cu" target="_blank">qualifiers.cu</a><br>
</span>@@ -1,21 +1,35 @@<br>
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s<br>
<span class=""> // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s<br>
+//<br>
</span>+// We run clang_cc1 with 'not' because source file contains<br>
+// intentional errors. CC1 failure is expected and must be ignored<br>
+// here. We're interested in what ends up in AST and that's what<br>
+// FileCheck verifies.<br></blockquote><div><br>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)))<br> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<span class="im HOEnZb">+// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -ast-dump %s \<br>
+// RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-HOST<br>
</span><div class="HOEnZb"><div class="h5">+// RUN: not %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -ast-dump -fcuda-is-device %s \<br>
+// RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-DEVICE<br>
<br>
#include "Inputs/cuda.h"<br>
<br>
// Host (x86) supports TLS and device-side compilation should ignore<br>
// host variables. No errors in either case.<br>
int __thread host_tls_var;<br>
+// CHECK-ALL: host_tls_var 'int' tls<br>
<br>
#if defined(__CUDA_ARCH__)<br>
// NVPTX does not support TLS<br>
__device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}}<br>
+// CHECK-DEVICE: device_tls_var 'int' tls<br>
__shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}}<br>
+// CHECK-DEVICE: shared_tls_var 'int' tls<br>
#else<br>
// Device-side vars should not produce any errors during host-side<br>
// compilation.<br>
__device__ int __thread device_tls_var;<br>
+// CHECK-HOST: device_tls_var 'int' tls<br>
__shared__ int __thread shared_tls_var;<br>
+// CHECK-HOST: shared_tls_var 'int' tls<br>
#endif<br>
<br>
__global__ void g1(int x) {}<br>
<br>
EMAIL PREFERENCES<br>
<a href="http://reviews.llvm.org/settings/panel/emailpreferences/" target="_blank">http://reviews.llvm.org/settings/panel/emailpreferences/</a><br>
</div></div><br>_______________________________________________<br>
cfe-commits mailing list<br>
<a href="mailto:cfe-commits@cs.uiuc.edu">cfe-commits@cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits</a><br>
<br></blockquote></div><br></div></div>