[PATCH] D58463: [CUDA]Delayed diagnostics for the asm instructions.

Alexey Bataev via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Tue Feb 26 11:41:41 PST 2019


ABataev added a comment.

In D58463#1411039 <https://reviews.llvm.org/D58463#1411039>, @tra wrote:

> > Hi Artem, I think we can just delay emission of this warning to solve this problem.
>
> I'm not sure we can always tell whether the warning is real or if it's the consequence of failing to parse inline asm.
>
> E.g.:
>
>   namespace {
>   __host__ __device__ a() {
>     int prev;
>     __asm__ __volatile__("mov %0, 0" : "=a" (prev)::);
>     return prev;
>   }
>  
>   __host__ __device__ b() {
>     int prev;
>     return prev;
>   }
>  
>   } //namespace
>
>
> Ideally we should always emit uninitialized diagnostics for `b`, but never for `a` in both host and device compilation modes.
>  I think we may want to propagate assignment from the inline asm statement -- we may not know the meaning of the constraint, but we do know which argument gets used/modified by the asm statement. Perhaps we can construct a fake GCCAsmStmt but bail out before we attempt to validate the asm string.


But it is going to be emitted for b() if b() is really used on the host or on the device. For a() the warning is going to be emitted only if it is really used on device, otherwise it is not.
Instead, we can try to do what we did before: construct GCCAsmStmt object, just like you said. What option do you prefer?


Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D58463/new/

https://reviews.llvm.org/D58463





More information about the llvm-commits mailing list