Using OpenACC

[ImanHosseini]

I have a suggestion: using OpenACC. OpenACC would support all 3 of simd, multicore and gpu and it is very convenient. I see that there is already cuda implementation of the kernels: OpenACC has CUDA interop, so the code before and after those cuda kernels can also be offloaded to GPU with OpenACC.

eval_ker_expts_libin_simd64

As a demonstration, I looked at 'eval_ker_expts_libin_simd64': rewriting the openmp loop with openacc is as easy as: (file: https://gist.github.com/ImanHosseini/c3ffc350c25d746cf8d524e6ede94573)

#pragma acc kernels
  for (int i=1;i<=M;++i) {  // changing from i=1 to i=0 slows from 0.2s to 2.4s!!!! I don't understand - has to be a better way to control (assembly code?)
    xi = i/(FLT)M;        // dummy offset to make each rep different
    for (int j=0;j<w;++j){
      tmp_val = -1.0 + xi + iw*j;
      ans += exp(beta * sqrt(1.0 - c*tmp_val*tmp_val));
    }
  }

The nice thing about this, is that the same code can be compiled for simd, multicore or gpu, and all it took was just 1 "acc kernels" pragma.

nvc++ t_acc.cpp -acc=multicore -Mvect=simd,noaltcode -fast -o t_acc_NV -Minfo=accel
nvc++ t_acc.cpp -acc=host -Mvect=simd,noaltcode -fast -o t_acc_NV -Minfo=accel
nvc++ t_acc.cpp -acc=gpu -fast -o t_acc_NV -Minfo=accel

On my machine (RTX 3060 Ti, 14-core i9-9940X):

Compilation	Time (s)
REF: g++, OpenMP simd	0.244
nvc++, OpenACC simd	0.325
nvc++, OpenACC multicore	0.057
nvc++, OpenACC gpu	0.229

Fortran

As a 2nd example, I took a fortran example: nufft1d_demo and changed that to use fortran doconcurrent and openacc (for 'errcomp') i.e.:

DO CONCURRENT (k1 = 1:ms)
            sk(k1) = 48 * dcos(k1 * pi / ms)
      ENDDO

and:

DO CONCURRENT (k1 = -nj/2 : (nj-1)/2)
            j = k1 + nj/2 + 1
            xj(j) = pi * dcos(-pi * j / nj)
            cj(j) = dcmplx(dsin(pi * j / nj), dcos(pi * j / nj))
      ENDDO

and:

!$acc parallel loop reduction(+:ealg,salg)
      do k = 1, n
            ealg = ealg + cdabs(fk1(k) - fk0(k))**2
            salg = salg + cdabs(fk0(k))**2
      enddo
!$acc end parallel loop

Compiled as:

nvfortran nufft1d_demo_acc.f ../directft/dirft1d.f -o nufft1d_demo_acc libfinufft_static.a -stdpar=multicore -acc=multicore -Minfo=accel -Mcudalib=cufft -lstdc++ -lfftw3 -lfftw3_omp -lm 

Result:

Compilation	Time (s)
REF: gfortran	3.785
nvfortran, Stdpar+OpenACC multicore	0.563