efriedma-quic wrote:
> Whilst I am thankful for the feedback I think it is somewhat unfortunate that
> we could not have a shared discussion about this, since I think that there
> are some core misunderstandings that seem to recur, which makes forward
> progress one way or the other difficult.
We didn't really say much on the call itself; we just spent a minute while we
were going through controversial RFCs/PRs, to call this out as something that
needed attention.
If you think this topic would benefit from a meeting, we can organize one...
but maybe a 1-on-1 chat would be better to start with, just to make sure we're
on the same page.
> The front-end cannot generate accurate diagnostics for the actual interesting
> case where the target is abstract (AMDGCNSPIRV, or the generic target
> @jhuber6 uses in GPU libc, if we extend things in that direction), because
> there isn't enough information - we only know what the concrete target is,
> and hence what features are available, at a point in time which is sequenced
> after the front-end has finished processing (could be run-time JIT for
> SPIR-V, could be bit code linking in a completely different compilation for
> GPU libc etc.);
If you have a construct like the following:
```
if (__builtin_amdgcn_processor_is("gfx900"))) {
some_gfx9000_specific_intrinsic();
}
some_gfx9000_specific_intrinsic()
```
We can tell, statically, that the first call is correctly guarded by an if
statement: it's guaranteed it will never run on a non-gfx9000 processor. The
second call, on the other hand, is not. So we can add a frontend rule: the
first call is legal, the second is not. Obviously the error has false
positives, in the sense that we can't actually prove the second call is
incorrect at runtime... but that's fine, probably.
What I don't want is that we end up with, essentially, the same constraint, but
enforced by the backend.
> There is not watertight mechanism here in the presence of indirect function
> calls / pointers to function
Sure; we can't stop people from calling arbitrary pointers.
> We do have a theoretical problem with guaranteeing that non-matching code
> isn't emitted, because LLVM IR doesn't promise to leave a code sequence like
> this alone:
There are ways to solve this: for example, we can make the
llvm.compiler.supports produce a token, and staple that token onto the
intrinsics using a bundle. Making this work requires that IRGen knows which
intrinic calls are actually impacted...
I care less about exactly how we solve this because we can adjust the solution
later. Whatever we expose in the frontend is much harder to change later.
https://github.com/llvm/llvm-project/pull/134016
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
