[PATCH] D104505: [HIP] Defer operator overloading errors

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D104505#2830996 <https://reviews.llvm.org/D104505#2830996>, @tra wrote:

> I don't think we want to do this.
>
>   struct S {
>       S& operator <<(int x);
>   };
>   
>   __device__ void foo() {
>       S s;
>       s<<1;
>   }
>   
>   <source>:7:6: error: invalid operands to binary expression ('S' and 'int')
>       s<<1;
>       ~^ ~
>   <source>:2:8: note: candidate function not viable: call to __host__ function from __device__ function
>       S& operator <<(int x);
>
> https://godbolt.org/z/Maa6Ed94W
> I believe diagnostic issued by clang is completely valid here.
>
> This is the case where clang and NVCC fundamentally differ in their compilation approach. NVCC effectively does not see device-side code during host compilation. Clang does. The code above is wrong regardless of whether we're compiling for the host of for the device and therefore, I believe, the diagnostic is appropriate.
>
> If the operator is intended to be used on GPU, it should have appropriate attributes. If it's not, then it's an error. NVCC not diagnosing it is a deficiency in NVCC, IMO. This particular problem should be fixed in the source code.

nvcc does see and parse the device functions in host compilation, e.g.

  # cat a.cu
  struct S {
      S& operator <<(int x);
  };
  
  #if !__CUDA_ARCH__
  __device__ void foo() {
      S s;
      s<<1;
      s<<;
  }
  #endif
  
  # nvcc a.cu
  a.cu(9): error: expected an expression

If nvcc simply skips all device functions, it will not diagnose the syntax error in the above code. It just chooses not to diagnose overloading resolution related errors.

P.S. To reproduce the above diagnostics you need to execute nvcc directly instead of on goldbolt.org, since glodbolt.org only does device compilation for CUDA.


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

https://reviews.llvm.org/D104505