This is an archive of the discontinued LLVM Phabricator instance.

Consider calls from implict host device functions as valid in SemaCUDA
ClosedPublic

Authored by jpienaar on Dec 8 2014, 10:50 AM.

Details

Summary

In SemaCUDA all implicit functions were considered host device, this led to errors such as the following code snippet failing to compile:

struct Copyable {
  const Copyable& operator=(const Copyable& x) { return *this; }
};

struct Simple {
  Copyable b;
};

void foo() {
  Simple a, b;

  a = b;
}

Above the implicit copy assignment operator was inferred as host device but there was only a host assignment copy defined which is an error in device compilation mode.

Diff Detail

Event Timeline

jpienaar updated this revision to Diff 17041.Dec 8 2014, 10:50 AM
jpienaar retitled this revision from to Consider only implicit builtin functions as device host in SemaCUDA.
jpienaar updated this object.
jpienaar edited the test plan for this revision. (Show Details)
jpienaar added reviewers: rnk, eliben.
jpienaar added a subscriber: Unknown Object (MLST).
rnk added inline comments.Dec 8 2014, 12:32 PM
lib/Sema/SemaCUDA.cpp
54

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?

test/SemaCUDA/implicit-copy.cu
16

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.

jpienaar added inline comments.Dec 8 2014, 2:43 PM
lib/Sema/SemaCUDA.cpp
54

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;
           ^
test/SemaCUDA/implicit-copy.cu
16

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.

jpienaar updated this revision to Diff 17054.Dec 8 2014, 2:45 PM

Changed the inference to treat all implicit functions uniformly. Extended test to include copy constructor.

jpienaar updated this revision to Diff 17206.Dec 11 2014, 7:51 PM
jpienaar retitled this revision from Consider only implicit builtin functions as device host in SemaCUDA to Consider calls from implict host device functions as valid in SemaCUDA.
jpienaar updated this object.

Avoided adding an additional attribute.

rnk accepted this revision.Dec 16 2014, 11:39 AM
rnk edited edge metadata.

lgtm

This revision is now accepted and ready to land.Dec 16 2014, 11:39 AM
This revision was automatically updated to reflect the committed changes.