2014-09-26 07:59:08 +08:00
|
|
|
// RUN: %clang_cc1 -fsyntax-only -verify %s
|
|
|
|
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 1: host method called from device function
|
|
|
|
|
|
|
|
struct S1 {
|
2016-08-10 09:09:18 +08:00
|
|
|
void method() {} // expected-note {{'method' declared here}}
|
2014-09-26 07:59:08 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
__device__ void foo1(S1& s) {
|
|
|
|
s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
|
|
|
|
}
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 2: host method called from device function, for overloaded method
|
|
|
|
|
|
|
|
struct S2 {
|
|
|
|
void method(int) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
|
|
|
|
void method(float) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
|
|
|
|
};
|
|
|
|
|
|
|
|
__device__ void foo2(S2& s, int i, float f) {
|
|
|
|
s.method(f); // expected-error {{no matching member function}}
|
|
|
|
}
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 3: device method called from host function
|
|
|
|
|
|
|
|
struct S3 {
|
[CUDA] Emit deferred diagnostics during Sema rather than during codegen.
Summary:
Emitting deferred diagnostics during codegen was a hack. It did work,
but usability was poor, both for us as compiler devs and for users. We
don't codegen if there are any sema errors, so for users this meant that
they wouldn't see deferred errors if there were any non-deferred errors.
For devs, this meant that we had to carefully split up our tests so that
when we tested deferred errors, we didn't emit any non-deferred errors.
This change moves checking for deferred errors into Sema. See the big
comment in SemaCUDA.cpp for an overview of the idea.
This checking adds overhead to compilation, because we have to maintain
a partial call graph. As a result, this change makes deferred errors a
CUDA-only concept (whereas before they were a general concept). If
anyone else wants to use this framework for something other than CUDA,
we can generalize at that time.
This patch makes the minimal set of test changes -- after this lands,
I'll go back through and do a cleanup of the tests that we no longer
have to split up.
Reviewers: rnk
Subscribers: cfe-commits, rsmith, tra
Differential Revision: https://reviews.llvm.org/D25541
llvm-svn: 284158
2016-10-14 04:52:12 +08:00
|
|
|
__device__ void method() {} // expected-note {{'method' declared here}}
|
2014-09-26 07:59:08 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
void foo3(S3& s) {
|
|
|
|
s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
|
|
|
|
}
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 4: device method called from host&device function
|
|
|
|
|
|
|
|
struct S4 {
|
[CUDA] Emit deferred diagnostics during Sema rather than during codegen.
Summary:
Emitting deferred diagnostics during codegen was a hack. It did work,
but usability was poor, both for us as compiler devs and for users. We
don't codegen if there are any sema errors, so for users this meant that
they wouldn't see deferred errors if there were any non-deferred errors.
For devs, this meant that we had to carefully split up our tests so that
when we tested deferred errors, we didn't emit any non-deferred errors.
This change moves checking for deferred errors into Sema. See the big
comment in SemaCUDA.cpp for an overview of the idea.
This checking adds overhead to compilation, because we have to maintain
a partial call graph. As a result, this change makes deferred errors a
CUDA-only concept (whereas before they were a general concept). If
anyone else wants to use this framework for something other than CUDA,
we can generalize at that time.
This patch makes the minimal set of test changes -- after this lands,
I'll go back through and do a cleanup of the tests that we no longer
have to split up.
Reviewers: rnk
Subscribers: cfe-commits, rsmith, tra
Differential Revision: https://reviews.llvm.org/D25541
llvm-svn: 284158
2016-10-14 04:52:12 +08:00
|
|
|
__device__ void method() {} // expected-note {{'method' declared here}}
|
2014-09-26 07:59:08 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
__host__ __device__ void foo4(S4& s) {
|
[CUDA] Emit deferred diagnostics during Sema rather than during codegen.
Summary:
Emitting deferred diagnostics during codegen was a hack. It did work,
but usability was poor, both for us as compiler devs and for users. We
don't codegen if there are any sema errors, so for users this meant that
they wouldn't see deferred errors if there were any non-deferred errors.
For devs, this meant that we had to carefully split up our tests so that
when we tested deferred errors, we didn't emit any non-deferred errors.
This change moves checking for deferred errors into Sema. See the big
comment in SemaCUDA.cpp for an overview of the idea.
This checking adds overhead to compilation, because we have to maintain
a partial call graph. As a result, this change makes deferred errors a
CUDA-only concept (whereas before they were a general concept). If
anyone else wants to use this framework for something other than CUDA,
we can generalize at that time.
This patch makes the minimal set of test changes -- after this lands,
I'll go back through and do a cleanup of the tests that we no longer
have to split up.
Reviewers: rnk
Subscribers: cfe-commits, rsmith, tra
Differential Revision: https://reviews.llvm.org/D25541
llvm-svn: 284158
2016-10-14 04:52:12 +08:00
|
|
|
s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
|
2014-09-26 07:59:08 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 5: overloaded operators
|
|
|
|
|
|
|
|
struct S5 {
|
|
|
|
S5() {}
|
|
|
|
S5& operator=(const S5&) {return *this;} // expected-note {{candidate function not viable}}
|
|
|
|
};
|
|
|
|
|
|
|
|
__device__ void foo5(S5& s, S5& t) {
|
|
|
|
s = t; // expected-error {{no viable overloaded '='}}
|
|
|
|
}
|
|
|
|
|
|
|
|
//------------------------------------------------------------------------------
|
|
|
|
// Test 6: call method through pointer
|
|
|
|
|
|
|
|
struct S6 {
|
2016-08-10 09:09:18 +08:00
|
|
|
void method() {} // expected-note {{'method' declared here}};
|
2014-09-26 07:59:08 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
__device__ void foo6(S6* s) {
|
|
|
|
s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
|
|
|
|
}
|