[PATCH] D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions

Thu Apr 22 10:19:11 PDT 2021

tra accepted this revision.
tra added a comment.
This revision is now accepted and ready to land.

In D100124#2707731 <https://reviews.llvm.org/D100124#2707731>, @steffenlarsen wrote:

>> Do you know if any existing code already uses the __nvvm_* builtins for cp.async? In other words, does nvcc provide them already or is it something we're free to name as we wish? I do not see any relevant intrinsics mentioned in NVVM IR spec: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html and I don't think NVCC's builtins are publicly documented anywhere.
>
> I don't know of any yet. We will be using these in the relatively near future, but we can still change them no problem. However, the intrinsic and builtin naming for NVVM and NVPTX seems a bit inconsistent so it may be a long discussion (or maybe not.)

LLVM uses different intrinsic names, mostly for historic reasons -- NVVM implements their own without upstreaming them back to LLVM and meanwhile LLVM grew its own set. So far I haven't seen any practical cases where that might've been an issue. I think NVIDIA folks popped up on a review *once* when they thought that an intrinsic we were about to introduce might've clashed with one of theirs, but they prompty disappeared when it turned out not to be the case.  The bottom line is that effectively intrinsic names in LLVM and NVVM are independent, though we should take care not to introduce identically named ones with different parameters or functionality.

Clang builtins are a bit different. Clang needs to compile CUDA headers and those  do use __nvvm builtinns, so clang must also provide those. NVIDIA does not document NVCC's compiler builtins, so if they are not used in CUDA headers, we have no idea whether relevant ones already exist. It would be great to stay in sync and make end-users code more portable across clang/NVCC, but there's not much we can do about that at the moment. The risk there is that if NVCC eventually introduces a builtin with the name we've used, but with a different arguments or functionality, that would be a bit of an annoyance for the users.

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

https://reviews.llvm.org/D100124