ldc nvvm GPU intrinsics good news

[Bruce Carneal]

On Friday, 5 March 2021 at 11:28:31 UTC, kinke wrote:
> On Friday, 5 March 2021 at 01:59:35 UTC, Bruce Carneal wrote:
[...]
>
> The dcompute parts in the compiler could definitely be improved 
> (but it's nowhere on my priorities list). What I tried to get 
> at by linking my comment [1] is that you can generally target 
> NVPTX/AMDGCN directly, without any @compute stuff, and in that 
> case importing and using intrinsics from that .di works without 
> any issues.
>
> [1] 
> https://github.com/ldc-developers/ldc/pull/3411#issuecomment-619385886

Okay.  I'll do the gccbuiltins_nvvm.di copy/edit for use in 
dcompute as part of my build script.

A copy/edit at build time will give access to many of the nvvm 
intrinsics and Johan's irEx suggestion can be used for access to 
others.

That leaves the intrinsics in the .td file that have llvm_i1_ty 
parameters and gcc builtin names (like vote/ballot).  IIUC these 
would all be handled by the gen_gccbuiltins.cpp program if it 
knew how to deal with llvm_i1_ty parameters.

Unfortunately, the "obvious" pragma(LDC_intrinsic, ...) forms 
that I've tried for llvm_i1_ty gccbuiltin intrinsics have not 
worked, at least not for the vote/ballot test cases.  I'll keep 
investigating.  Clang and nvcc are both based on LLVM and deal 
with these somehow so I'll start looking there.

I expect that you're very busy but if you already know or suspect 
what the llvm_i1_ty problem/fix is, a brief sketch here would be 
much appreciated.  Regardless, thanks again for your efforts.  A 
working LDC/dcompute is *much* nicer than C++/CUDA.