[PATCH] D47201: [CUDA] Implement nv_weak attribute for functions

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Tue Jun 5 16:15:19 PDT 2018


tra added a comment.

I've experimented a bit and I think that we may not need this patch at all.
As far as I can tell, nv_weak is only applicable to __device__ functions. It's ignored for __global__ kernels and is apparently forbidden for data.
For __device__ functions nvcc produces .weak attribute in PTX.

Using plain old __attribute__((weak)) does exactly the same. Considering that nv_weak is only used inside CUDA SDK headers, substituting weak in place of nv_weak will result in correct PTX, which is all we really need. I don't see much benefit turning it into full blown attribute just to mimic an internal CUDA implementation detail we don't care all that much about.

Now, replacing it in CUDA headers for all CUDA versions we support may be tricky. Let me give it a try. I'll send a patch, if I manage to make it work.


Repository:
  rC Clang

https://reviews.llvm.org/D47201





More information about the cfe-commits mailing list