================
Comment at: lib/Sema/SemaCUDA.cpp:54
@@ -53,4 +53,3 @@
return CFT_Host;
- } else if (D->isImplicit()) {
- // Some implicit declarations (like intrinsic functions) are not marked.
- // Set the most lenient target on them for maximal flexibility.
+ } else if (D->isImplicit() && D->getBuiltinID()) {
+ // Intrinsic functions cannot (at present) be marked as __device__ so set
----------------
rnk wrote:
> Builtins are for things like __builtin_popcount(), but your test case is
> looking at implicitly defined C++ special members, right? What was happening
> on the test case you have previously?
Yes, although that might be too restrictive.
I was getting the following errors:
```
error: object of type 'Simple' cannot be assigned because its copy assignment
operator is implicitly deleted
a = b;
^
note: copy assignment operator of 'Simple' is implicitly deleted because field
'b' has no copy assignment operator
Copyable b;
^
```
================
Comment at: test/SemaCUDA/implicit-copy.cu:16
@@ +15,2 @@
+ a = b;
+}
----------------
rnk wrote:
> This seems like another interesting test case:
>
> struct Copyable {
> __device__ const Copyable& operator=(const Copyable& x) { return *this; }
> };
> struct Simple {
> Copyable b;
> };
> void foo(Simple &a, Simple &b) {
> a = b;
> }
>
> Simple's implicit copy ctor should be rejected when the host is the target
> because it calls a device-only method from the host, right?
>
> I'm also curious what happens when someone uses FP math builtins like
> __builtin_cos that might have a reasonable lowering on a GPU.
Yes, and it does that. The one where there is a problem is with
```
__device__ void foo(Simple &a, Simple &b) {
a = b
}
```
There an error is flagged. In general, as CheckCUDATarget only looks at Caller
and Callee, will we run into a problem in cases such as below
```
Host function calls a implicit function which calls a Device function
```
Now, a host function calling an implicit function is fine, an implicit function
calling a device function is fine, but transitively we have a host function
calling a device function. But this all seem to be just used for the method
candidate set and errors are still reported if this happens. So it would seem
that we could allow implicits in here while still catching such erroneous cases.
I have created a diff where implicits are considered always OK, I'll append it.
http://reviews.llvm.org/D6565
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
