This is an archive of the discontinued LLVM Phabricator instance.

[CUDA][HIP] Always defer diagnostics for wrong-sided reference
ClosedPublic

Authored by yaxunl on Jul 15 2020, 11:21 AM.

Details

Summary

When a device function calls a host function or vice versa, this is wrong-sided
reference. Currently clang immediately diagnose it. This is different from nvcc
behavior, where it is diagnosed only if the function is really emitted:

void foo();
inline __device__ void bar() { foo(); }

https://godbolt.org/z/4G8P31

To me nvcc behavior makes more sense, since we do not need a diagnostic
unless we really need it. Current clang behavior causes diagnostics for wrong-sided
reference emitted twice, once in host compilation, once in device compilation,
which is unnecessary and cluttering the screen.

More importantly, current clang behavior causes false alarms for valid use cases:

__device__ void bar(int x);

template<class T>
__device__ void foo(T x) {
  bar(x);
}

struct A {
  int a;
  operator int() { return a; }
};

void foo(A a) {}

template<typename T>
__device__ __host__ void test(T t) {
  foo(t);
}

int main() {
  A a;
  test(a);
  return 0;
}

In this example, we only emit test<A> on host, it should compile.
However since test<A> is a host device function, it triggers
instantiation of foo<A> in device compilation. foo<A> calls
a host operator int() which causes immediate diagnostics.

In this case, user may never intend to use foo<A> on device side.
The instantiation of foo<A> is a side effect of instantiation of
host device function test<A>, therefore any diagnostics incurred
by foo<A> should be deferred.

This patch let clang always defer diagnostics for wrong-sided
reference.

Diff Detail

Event Timeline

yaxunl created this revision.Jul 15 2020, 11:21 AM

This is different from nvcc behavior, where it is diagnosed only if the function is really emitted:

That by itself is not a sufficient reason for relaxing the checks. Clang is stricter/more principled in diagnosing other questionalbe things nvcc lets through, even when it should not have.
NVCC only sees part of the source, so it's limited in what it can diagnose during particular compilation side. Clang can do better.

Current clang behavior causes diagnostics for wrong-sided reference emitted twice, once in host compilation, once in device compilation, which is unnecessary and cluttering the screen.

I don't think it's a problem. Clang will abort compilation if one of sub-compilation fails. We'll only see the error once.

More importantly, current clang behavior causes false alarms for valid use cases:

I'm not convinced that your example is valid. To me it looks like test should not have been a HD, given that it always uses a host function.

That said, I can see where it may be useful in case of test being a constexpr template function (for example somewhere in a standard library) parametrized by something else which may be host or device-only.
E.g. this may be a somewhat better example: https://godbolt.org/z/44vPv8

OK. I think it is useful, but I could also use a second opinion.
@rsmith , @jlebar do you have any thoughts on this?

tra and I talked offline and I...think this makes sense.

tra accepted this revision.Jul 16 2020, 3:03 PM

LGTM.

This revision is now accepted and ready to land.Jul 16 2020, 3:03 PM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptJul 17 2020, 4:52 AM