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

[Artem Belevich via Phabricator via cfe-commits]

tra added a comment.

>> 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.

Clang also emits the uninitialized warnings for `b` when it is not used -- as in the example above.
I'm OK with that as `b` is a valid function on both sides.

Suppressing uninitialized warning in this case would be wrong, IMO -- that would diverge from what clang would do if `b` didn't have `__host__ __device__` attributes.

> 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?

I think creating a GCCAsmStmt() is the right way to deal with this as it gives compiler the correct (well, as correct as we can at that point) info about the code, as opposed to giving compiler broken pieces and trying to suppress the fallout.


Repository:
  rL LLVM

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

https://reviews.llvm.org/D58463