1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
|
// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -std=c++17 -verify %s
#include "Inputs/cuda.h"
// Error, instantiated on device.
inline __host__ __device__ void hasInvalid() {
throw NULL;
// expected-error@-1 2{{cannot use 'throw' in __host__ __device__ function}}
}
inline __host__ __device__ void hasInvalid2() {
throw NULL;
// expected-error@-1 2{{cannot use 'throw' in __host__ __device__ function}}
}
inline __host__ __device__ void hasInvalidDiscarded() {
// This is only used in the discarded statements below, so this should not diagnose.
throw NULL;
}
static __device__ void use0() {
hasInvalid(); // expected-note {{called by 'use0'}}
hasInvalid(); // expected-note {{called by 'use0'}}
if constexpr (true) {
hasInvalid2(); // expected-note {{called by 'use0'}}
} else {
hasInvalidDiscarded();
}
if constexpr (false) {
hasInvalidDiscarded();
} else {
hasInvalid2(); // expected-note {{called by 'use0'}}
}
if constexpr (false) {
hasInvalidDiscarded();
}
}
// To avoid excessive diagnostic messages, deferred diagnostics are only
// emitted the first time a function is called.
static __device__ void use1() {
use0(); // expected-note 4{{called by 'use1'}}
use0();
}
static __device__ void use2() {
use1(); // expected-note 4{{called by 'use2'}}
use1();
}
static __device__ void use3() {
use2(); // expected-note 4{{called by 'use3'}}
use2();
}
__global__ void use4() {
use3(); // expected-note 4{{called by 'use4'}}
use3();
}
|
