This revision was automatically updated to reflect the committed changes.
Closed by commit rL284553: [CUDA] Rework tests now that we emit deferred
diagnostics during sema. Test… (authored by jlebar).
Changed prior to commit:
https://reviews.llvm.org/D25755?vs=75085&id=75099#toc
Repository:
rL LLVM
https://reviews.llvm.org/D25755
Files:
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-host-device.cu
cfe/trunk/test/SemaCUDA/exceptions.cu
cfe/trunk/test/SemaCUDA/function-overload-hd.cu
cfe/trunk/test/SemaCUDA/function-overload.cu
cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
cfe/trunk/test/SemaCUDA/static-vars-hd.cu
cfe/trunk/test/SemaCUDA/vla-host-device.cu
cfe/trunk/test/SemaCUDA/vla.cu
Index: cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu
===================================================================
--- cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu
+++ cfe/trunk/test/PCH/pragma-cuda-force-host-device.cu
@@ -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
Index: cfe/trunk/test/Parser/cuda-force-host-device-templates.cu
===================================================================
--- cfe/trunk/test/Parser/cuda-force-host-device-templates.cu
+++ cfe/trunk/test/Parser/cuda-force-host-device-templates.cu
@@ -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}}
Index: cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
+++ cfe/trunk/test/SemaCUDA/implicit-device-lambda.cu
@@ -76,6 +76,26 @@
f4();
}
+__host__ __device__ void hd_fn() {
+ auto f1 = [&] {};
+ f1(); // implicitly __host__ __device__
+
+ auto f2 = [&] __device__ {};
+ f2();
+#ifndef __CUDA_ARCH__
+ // expected-error@-2 {{reference to __device__ function}}
+#endif
+
+ auto f3 = [&] __host__ {};
+ f3();
+#ifdef __CUDA_ARCH__
+ // expected-error@-2 {{reference to __host__ function}}
+#endif
+
+ auto f4 = [&] __host__ __device__ {};
+ f4();
+}
+
// The special treatment above only applies to lambdas.
__device__ void foo() {
struct X {
Index: cfe/trunk/test/SemaCUDA/vla.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vla.cu
+++ cfe/trunk/test/SemaCUDA/vla.cu
@@ -10,3 +10,16 @@
__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@-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); }
Index: cfe/trunk/test/SemaCUDA/function-overload.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu
+++ cfe/trunk/test/SemaCUDA/function-overload.cu
@@ -193,12 +193,22 @@
CurrentFnPtr fp_cdh = cdh;
CurrentReturnTy ret_cdh = cdh();
+ GlobalFnPtr fp_g = g;
+#if defined(__CUDA_ARCH__)
+ // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
+
g();
#if defined (__CUDA_ARCH__)
// expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
#else
// expected-error@-4 {{call to global function g not configured}}
#endif
+
+ g<<<0,0>>>();
+#if defined(__CUDA_ARCH__)
+ // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
}
// Test for address of overloaded function resolution in the global context.
Index: cfe/trunk/test/SemaCUDA/device-var-init.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu
@@ -213,3 +213,15 @@
static int v;
// expected-error@-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@-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(); }
Index: cfe/trunk/test/SemaCUDA/exceptions.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/exceptions.cu
+++ cfe/trunk/test/SemaCUDA/exceptions.cu
@@ -19,3 +19,34 @@
try {} catch(void*) {}
// expected-error@-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@-3 {{cannot use 'throw' in __host__ __device__ function}}
+ // expected-error@-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@-3 {{cannot use 'throw' in __host__ __device__ function}}
+ // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
+#endif
+}
+
+__device__ void call_hd3() { hd3(); }
Index: cfe/trunk/test/SemaCUDA/static-vars-hd.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/static-vars-hd.cu
+++ cfe/trunk/test/SemaCUDA/static-vars-hd.cu
@@ -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@-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(); }
Index: cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
+++ cfe/trunk/test/SemaCUDA/implicit-device-lambda-hd.cu
@@ -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@-2 {{reference to __device__ function}}
-#endif
-
- auto f3 = [&] __host__ {};
- f3();
-#ifndef HOST
- // expected-error@-2 {{reference to __host__ function}}
-#endif
-
- auto f4 = [&] __host__ __device__ {};
- f4();
-}
-
-
Index: cfe/trunk/test/SemaCUDA/vla-host-device.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vla-host-device.cu
+++ cfe/trunk/test/SemaCUDA/vla-host-device.cu
@@ -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@-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); }
Index: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
+++ cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
@@ -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@-3 {{cannot use 'throw' in __host__ __device__ function}}
- // expected-error@-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@-3 {{cannot use 'throw' in __host__ __device__ function}}
- // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
-#endif
-}
-__device__ void call_hd3() { hd3(); }
Index: cfe/trunk/test/SemaCUDA/function-overload-hd.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/function-overload-hd.cu
+++ cfe/trunk/test/SemaCUDA/function-overload-hd.cu
@@ -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@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif
- g<<<0,0>>>();
-#if defined(__CUDA_ARCH__)
- // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif // __CUDA_ARCH__
-}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
