r284553 - [CUDA] Rework tests now that we emit deferred diagnostics during sema. Test-only change.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 18 17:06:49 PDT 2016
Author: jlebar
Date: Tue Oct 18 19:06:49 2016
New Revision: 284553
URL: http://llvm.org/viewvc/llvm-project?rev=284553&view=rev
Log:
[CUDA] Rework tests now that we emit deferred diagnostics during sema. Test-only change.
Summary:
Previously we had to split out a lot of our tests into a test that
checked only immediate errors and a test that checked only deferred
errors. This was because, if you emitted any immediate errors, we
wouldn't run codegen, where the deferred errors were emitted.
We've fixed this, and now emit deferred errors during sema. This lets
us merge a bunch of tests, and lets us convert some other tests to
-fsyntax-only.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D25755
Removed:
cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
cfe/trunk/test/SemaCUDA/function-overload-hd.cu
cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
cfe/trunk/test/SemaCUDA/static-vars-hd.cu
cfe/trunk/test/SemaCUDA/vla-host-device.cu
Modified:
cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu
cfe/trunk/test/Parser/cuda-force-host-device-templates.cu
cfe/trunk/test/SemaCUDA/device-var-init.cu
cfe/trunk/test/SemaCUDA/exceptions.cu
cfe/trunk/test/SemaCUDA/function-overload.cu
cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
cfe/trunk/test/SemaCUDA/vla.cu
Modified: cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu (original)
+++ cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu Tue Oct 18 19:06:49 2016
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -emit-pch %s -o %t
-// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -S -o /dev/null %s
+// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -fsyntax-only %s
#ifndef HEADER
#define HEADER
Modified: cfe/trunk/test/Parser/cuda-force-host-device-templates.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Parser/cuda-force-host-device-templates.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/Parser/cuda-force-host-device-templates.cu (original)
+++ cfe/trunk/test/Parser/cuda-force-host-device-templates.cu Tue Oct 18 19:06:49 2016
@@ -1,8 +1,7 @@
-// RUN: %clang_cc1 -std=c++14 -S -verify -fcuda-is-device %s -o /dev/null
+// RUN: %clang_cc1 -std=c++14 -fsyntax-only -verify -fcuda-is-device %s
// Check how the force_cuda_host_device pragma interacts with template
-// instantiations. The errors here are emitted at codegen, so we can't do
-// -fsyntax-only.
+// instantiations.
template <typename T>
auto foo() { // expected-note {{declared here}}
Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Tue Oct 18 19:06:49 2016
@@ -213,3 +213,15 @@ __device__ void df_sema() {
static int v;
// expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
}
+
+__host__ __device__ void hd_sema() {
+ static int x = 42;
+#ifdef __CUDA_ARCH__
+ // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+#endif
+}
+
+inline __host__ __device__ void hd_emitted_host_only() {
+ static int x = 42; // no error on device because this is never codegen'ed there.
+}
+void call_hd_emitted_host_only() { hd_emitted_host_only(); }
Removed: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions-host-device.cu?rev=284552&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/exceptions-host-device.cu (original)
+++ cfe/trunk/test/SemaCUDA/exceptions-host-device.cu (removed)
@@ -1,38 +0,0 @@
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null
-// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null
-
-#include "Inputs/cuda.h"
-
-// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
-// function if and only if it's codegen'ed for device.
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void hd1() {
- throw NULL;
- try {} catch(void*) {}
-#ifndef HOST
- // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function}}
- // expected-error at -3 {{cannot use 'try' in __host__ __device__ function}}
-#endif
-}
-
-// No error, never instantiated on device.
-inline __host__ __device__ void hd2() {
- throw NULL;
- try {} catch(void*) {}
-}
-void call_hd2() { hd2(); }
-
-// Error, instantiated on device.
-inline __host__ __device__ void hd3() {
- throw NULL;
- try {} catch(void*) {}
-#ifndef HOST
- // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function}}
- // expected-error at -3 {{cannot use 'try' in __host__ __device__ function}}
-#endif
-}
-__device__ void call_hd3() { hd3(); }
Modified: cfe/trunk/test/SemaCUDA/exceptions.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/exceptions.cu (original)
+++ cfe/trunk/test/SemaCUDA/exceptions.cu Tue Oct 18 19:06:49 2016
@@ -19,3 +19,34 @@ __global__ void kernel() {
try {} catch(void*) {}
// expected-error at -1 {{cannot use 'try' in __global__ function}}
}
+
+// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
+// function if and only if it's codegen'ed for device.
+
+__host__ __device__ void hd1() {
+ throw NULL;
+ try {} catch(void*) {}
+#ifdef __CUDA_ARCH__
+ // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function}}
+ // expected-error at -3 {{cannot use 'try' in __host__ __device__ function}}
+#endif
+}
+
+// No error, never instantiated on device.
+inline __host__ __device__ void hd2() {
+ throw NULL;
+ try {} catch(void*) {}
+}
+void call_hd2() { hd2(); }
+
+// Error, instantiated on device.
+inline __host__ __device__ void hd3() {
+ throw NULL;
+ try {} catch(void*) {}
+#ifdef __CUDA_ARCH__
+ // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function}}
+ // expected-error at -3 {{cannot use 'try' in __host__ __device__ function}}
+#endif
+}
+
+__device__ void call_hd3() { hd3(); }
Removed: cfe/trunk/test/SemaCUDA/function-overload-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload-hd.cu?rev=284552&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload-hd.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload-hd.cu (removed)
@@ -1,31 +0,0 @@
-// REQUIRES: x86-registered-target
-// REQUIRES: nvptx-registered-target
-
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -S -o /dev/null -verify \
-// RUN: -verify-ignore-unexpected=note %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -S -o /dev/null -fcuda-is-device \
-// RUN: -verify -verify-ignore-unexpected=note %s
-
-#include "Inputs/cuda.h"
-
-// FIXME: Merge into function-overload.cu once deferred errors can be emitted
-// when non-deferred errors are present.
-
-#if !defined(__CUDA_ARCH__)
-//expected-no-diagnostics
-#endif
-
-typedef void (*GlobalFnPtr)(); // __global__ functions must return void.
-
-__global__ void g() {}
-
-__host__ __device__ void hd() {
- GlobalFnPtr fp_g = g;
-#if defined(__CUDA_ARCH__)
- // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif
- g<<<0,0>>>();
-#if defined(__CUDA_ARCH__)
- // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif // __CUDA_ARCH__
-}
Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Tue Oct 18 19:06:49 2016
@@ -193,12 +193,22 @@ __host__ __device__ void hostdevicef() {
CurrentFnPtr fp_cdh = cdh;
CurrentReturnTy ret_cdh = cdh();
+ GlobalFnPtr fp_g = g;
+#if defined(__CUDA_ARCH__)
+ // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
+
g();
#if defined (__CUDA_ARCH__)
// expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
#else
// expected-error at -4 {{call to global function g not configured}}
#endif
+
+ g<<<0,0>>>();
+#if defined(__CUDA_ARCH__)
+ // expected-error at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
}
// Test for address of overloaded function resolution in the global context.
Removed: cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu?rev=284552&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu (removed)
@@ -1,27 +0,0 @@
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -verify-ignore-unexpected=note \
-// RUN: -S -o /dev/null %s
-// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=note \
-// RUN: -DHOST -S -o /dev/null %s
-#include "Inputs/cuda.h"
-
-__host__ __device__ void hd_fn() {
- auto f1 = [&] {};
- f1(); // implicitly __host__ __device__
-
- auto f2 = [&] __device__ {};
- f2();
-#ifdef HOST
- // expected-error at -2 {{reference to __device__ function}}
-#endif
-
- auto f3 = [&] __host__ {};
- f3();
-#ifndef HOST
- // expected-error at -2 {{reference to __host__ function}}
-#endif
-
- auto f4 = [&] __host__ __device__ {};
- f4();
-}
-
-
Modified: cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu Tue Oct 18 19:06:49 2016
@@ -76,6 +76,26 @@ __host__ void host_fn() {
f4();
}
+__host__ __device__ void hd_fn() {
+ auto f1 = [&] {};
+ f1(); // implicitly __host__ __device__
+
+ auto f2 = [&] __device__ {};
+ f2();
+#ifndef __CUDA_ARCH__
+ // expected-error at -2 {{reference to __device__ function}}
+#endif
+
+ auto f3 = [&] __host__ {};
+ f3();
+#ifdef __CUDA_ARCH__
+ // expected-error at -2 {{reference to __host__ function}}
+#endif
+
+ auto f4 = [&] __host__ __device__ {};
+ f4();
+}
+
// The special treatment above only applies to lambdas.
__device__ void foo() {
struct X {
Removed: cfe/trunk/test/SemaCUDA/static-vars-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/static-vars-hd.cu?rev=284552&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/static-vars-hd.cu (original)
+++ cfe/trunk/test/SemaCUDA/static-vars-hd.cu (removed)
@@ -1,20 +0,0 @@
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s
-// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s
-
-#include "Inputs/cuda.h"
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void f() {
- static int x = 42;
-#ifndef HOST
- // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
-#endif
-}
-
-inline __host__ __device__ void g() {
- static int x = 42; // no error on device because this is never codegen'ed there.
-}
-void call_g() { g(); }
Removed: cfe/trunk/test/SemaCUDA/vla-host-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vla-host-device.cu?rev=284552&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/vla-host-device.cu (original)
+++ cfe/trunk/test/SemaCUDA/vla-host-device.cu (removed)
@@ -1,21 +0,0 @@
-// RUN: %clang_cc1 -fcuda-is-device -verify -S %s -o /dev/null
-// RUN: %clang_cc1 -verify -DHOST %s -S -o /dev/null
-
-#include "Inputs/cuda.h"
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void hd(int n) {
- int x[n];
-#ifndef HOST
- // expected-error at -2 {{cannot use variable-length arrays in __host__ __device__ functions}}
-#endif
-}
-
-// No error because never codegen'ed for device.
-__host__ __device__ inline void hd_inline(int n) {
- int x[n];
-}
-void call_hd_inline() { hd_inline(42); }
Modified: cfe/trunk/test/SemaCUDA/vla.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vla.cu?rev=284553&r1=284552&r2=284553&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/vla.cu (original)
+++ cfe/trunk/test/SemaCUDA/vla.cu Tue Oct 18 19:06:49 2016
@@ -10,3 +10,16 @@ void host(int n) {
__device__ void device(int n) {
int x[n]; // expected-error {{cannot use variable-length arrays in __device__ functions}}
}
+
+__host__ __device__ void hd(int n) {
+ int x[n];
+#ifdef __CUDA_ARCH__
+ // expected-error at -2 {{cannot use variable-length arrays in __host__ __device__ functions}}
+#endif
+}
+
+// No error because never codegen'ed for device.
+__host__ __device__ inline void hd_inline(int n) {
+ int x[n];
+}
+void call_hd_inline() { hd_inline(42); }
More information about the cfe-commits
mailing list