[PATCH] D70172: [CUDA][HIP] Fix assertion due to dtor check on windows

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D70172#1810665 <https://reviews.llvm.org/D70172#1810665>, @rsmith wrote:

> This doesn't look quite right to me. I don't think we should treat the `delete this;` for a destructor as being emitted-for-device in any translation unit in which the vtable is marked used. (For example, if in your testcase `MSEmitDeletingDtor::CFileStream::CFileStream()` were a `__host__` function, I think you'd still diagnose, but presumably shouldn't do so, because the vtable -- and therefore `CFileStream::operator delete` -- is never referenced / emitted for the device.) Instead, I think we should treat the `delete this;` as being emitted in any translation unit in which the vtable itself is emitted-for-device. Presumably, this means you will need to model / track usage of the vtable itself in your "call graph".


A user declared ctor/dtor by default is `__host__`.

Let's consider this testcase:

  static __device__ __host__ void f(__m256i *p) {
    __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
                   : "r0"); // MS-error{{unknown register name 'r0' in asm}}
  }
  struct CFileStream {
    void operator delete(void *p) {
      f(0);  // MS-note{{called by 'operator delete'}}
    }
    CFileStream();
    virtual ~CFileStream();  // MS-note{{called by '~CFileStream'}}
  };
  
  struct CMultiFileStream {
    CFileStream m_fileStream;
    ~CMultiFileStream();
  };
  
  // This causes vtable emitted so that deleting dtor is emitted for MS.
  CFileStream::CFileStream() {}

In host compilation, vtbl is emitted, since it causes dtor emitted, whereas dtor calls f(), therefore the diagnostic msg is emitted.

In device compilation, vtbl is not emitted, therefore dtor is not emitted, and the diagnostic msg in f() is not emitted.

We only need an entity in call graph if that entity can be called by other entities. Here vtbl is always at the top level of the 'call graph'. Therefore it is not needed to be in the call graph.


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

https://reviews.llvm.org/D70172